ó
    “~i\‹ ã                  óº   • S SK Jr  S SKrS SKJrJrJrJrJrJ	r	J
r
  S SKrS SKJr  SSKJr  SSKJr  \" S	5      r\" S
5      r " S S\5      r " S S\	\   5      rg)é    )ÚannotationsN)ÚListÚOptionalÚSequenceÚTupleÚTypeVarÚGenericÚType)Údriveré   )Úiré   )ÚcoreÚTÚTensorTyc                  ó(   ^ • \ rS rSrU 4S jrSrU =r$ )ÚIncompatibleTypeErrorImplé   c                óØ   >• Xl         X l        SU R                   R                  5       -   S-   U R                  R                  5       -   U l        [        [
        U ]  U R                  5        g )Nzinvalid operands of type ú and )Útype_aÚtype_bÚ__repr__ÚmessageÚsuperr   Ú__init__)Úselfr   r   Ú	__class__s      €ÚU/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/language/semantic.pyr   Ú"IncompatibleTypeErrorImpl.__init__   sT   ø€ ØŒØŒØ2°T·[±[×5IÑ5IÓ5KÑKÈgÑUÐX\×XcÑXc×XlÑXlÓXnÑnˆŒÜÔ'¨Ñ7¸¿¹ÕEó    )r   r   r   )Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__r   Ú__static_attributes__Ú__classcell__)r   s   @r   r   r      s   ø† ÷Fó Fr!   r   c                  óD  • \ rS rSr% \R
                  rS\S'   \rS\S'   S rSS jr	SS jr
S€S	 jr    S�S
 jrS‚SƒS jjrS„S jr  S… S†S jjrS‡S jr    SˆS jr    SˆS jr    SˆS jrS‰S jrS‰S jrSŠS jrS‰S jrS‹S jrS‹S jrSŒS jrS�S jrSŽS jrSŽS jrSŽS jrSŽS jrSŽS jr S�S jr!SŽS  jr"SŽS! jr#SŽS" jr$S�S# jr%S�S$ jr&S�S% jr'S‘S& jr(SŽS' jr)SŽS( jr*SŽS) jr+SŽS* jr,SŽS+ jr-SŽS, jr.S-S..S’S/ jjr/S“S0 jr0S“S1 jr1S”S2 jr2S•S3 jr3S–S4 jr4S—S5 jr5S˜S6 jr6S™S7 jr7SšS8 jr8S›S9 jr9SœS: jr:S�S; jr;SžS< jr<SŸS= jr=S S> jr>S¡S¢S? jjr?S@ r@SA rASB rBSC rCSD rDSE rESF rFSG rGSH rH          S£SI jrI    S¤SJ jrJS¥SK jrKS¦SL jrLS¦SM jrMSN rNSO rOS¦SP jrPS¦SQ jrQS¦SR jrRS¦SS jrSS¦ST jrTS§SU jrUS–SV jrVSW rWSX rX    S¨SY jrYS©SZ jrZ    SªS[ jr[S«S\ jr\S¬S] jr]S¬S^ jr^S¬S_ jr_S¬S` jr`S¬Sa jraS¬Sb jrbS¬Sc jrcSd rd      S­Se jreS®Sf jrfS¯Sg jrgSh rh                S°Si jriS±Sj jrjSk rkS²Sl jrl    S³Sm jrmS´Sn jrnSo ro  SµSp jrpS¶Sq jrqS·Sr jrrS·Ss jrsS·St jrtS¸Su jruS¹Sv jrvSºSw jrwS¸Sx jrxSy ryS‚Sz jrzS»S{ jr{S»S| jr| S¼     S½S} jjr}S~r~g-)¾ÚTritonSemanticé   zType[TensorTy]Útensorz
ir.builderÚbuilderc                ó   • Xl         g ©N©r,   )r   r,   s     r   r   ÚTritonSemantic.__init__   s   € Ø�r!   c                óœ   • US;  a  [        SU 35      eU R                  U R                  R                  U5      [        R
                  5      $ )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_idÚTritonSemantic.program_id&   sA   € Ø�yÓ ÜÐJÈ4È&ÐQÓRÐRØ�{‰{˜4Ÿ<™<×=Ñ=¸dÓCÄRÇXÁXÓNÐNr!   c                óœ   • US;  a  [        SU 35      eU R                  U R                  R                  U5      [        R
                  5      $ )Nr2   z-num_programs axis must be 0, 1, or 2 but got )r3   r+   r,   Úcreate_get_num_programsr5   r6   r7   s     r   Únum_programsÚTritonSemantic.num_programs+   sA   € Ø�yÓ ÜÐLÈTÈFÐSÓTÐTØ�{‰{˜4Ÿ<™<×?Ñ?ÀÓEÄrÇxÁxÓPÐPr!   c                ód  • UR                   nUR                   nUR                  nUR                  nXV:X  a	  X4:”  a  U$ U$ U[        R                  R                  R
                  :X  a	  X4:¼  a  U$ U$ U[        R                  R                  R
                  :X  a	  XC:¼  a  U$ U$ [        SU SU 35      e)Nzunexpected signedness r   )Úint_bitwidthÚint_signednessr5   ÚdtypeÚ
SIGNEDNESSÚUNSIGNEDÚ	TypeError)r   Úa_tyÚb_tyÚa_rankÚb_rankÚa_snÚb_sns          r   Úinteger_promote_implÚ#TritonSemantic.integer_promote_impl4   s­   € Ø×"Ñ"ˆØ×"Ñ"ˆØ×"Ñ"ˆØ×"Ñ"ˆð ‹<Ø!›?�4Ð4°Ð4Ø”R—X‘X×(Ñ(×1Ñ1Ó1Ø!Ó+�4Ð5°Ð5Ø”R—X‘X×(Ñ(×1Ñ1Ó1Ø!Ó+�4Ð5°Ð5ÜÐ0°°°e¸D¸6ÐBÓCÐCr!   c                ór  • X$:w  a‚  U(       a  X4OX14u  pgUR                  5       R                  UR                  5       R                  ::  a=  U(       a4  U[        R                  [        R                  4;   a  [        R
                  $ U$ UR                  5       (       d  UR                  5       (       a  [        R                  $ UR                  5       (       d  UR                  5       (       a  [        R
                  $ UR                  5       (       d  UR                  5       (       a'  U(       a  [        R
                  $ [        R                  $ UR                  5       (       a<  UR                  5       (       a'  U(       a  [        R
                  $ [        R                  $ UR                  5       (       d  UR                  5       (       a  [        R
                  $ UR                  5       (       a,  UR                  5       (       a  X:X  a  U$ [        R                  $ UR                  5       (       a  UR                  5       (       d  [        SU SU 35      eU(       aM  UR                  UR                  :w  a3  [        SUR                  5       -   S-   UR                  5       -   S-   5      eU R!                  X5      $ )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Úvaluer5   Úfloat16Úbfloat16Úfloat32Úis_fp64Úfloat64Úis_fp32Úis_fp16Úis_bf16Úis_fp8Úis_intrE   rA   r   rL   )r   rF   Úa_is_scalarrG   Úb_is_scalarÚ
div_or_modÚ	scalar_tyÚ	tensor_tys           r   Úcomputation_type_implÚ$TritonSemantic.computation_type_implC   sÜ  € ð
 Ó%Þ3> D¡<ÀTÀLÑ ˆIØ�~‰~Ó×%Ñ%¨¯©Ó)9×)?Ñ)?Ó?æ 9´·±¼R¿[¹[Ð0IÓ#IÜŸ:™:Ð%Ø Ð ð �<‰<�>‰>˜TŸ\™\Ÿ^™^Ü—:‘:Ðð �<‰<�>‰>˜TŸ\™\Ÿ^™^Ü—:‘:Ðð �<‰<�>‰>˜TŸ\™\Ÿ^™^ÞÜ—z‘zÐ!ä—z‘zÐ!à�<‰<�>‰>˜dŸl™lŸn™nÞÜ—z‘zÐ!ä—{‘{Ð"Ø�<‰<�>‰>˜TŸ\™\Ÿ^™^Ü—:‘:Ðà�;‰;�=‰=˜TŸ[™[Ÿ]™]Ø›<�4Ð7¬R¯Z©ZÐ7Ø�{‰{�}‰} D§K¡K§M¡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ð ×(Ñ(¨Ó4Ð4r!   c                ó  • [        U[        5      (       a9  U R                  U R                  R	                  U5      [
        R                  5      $ [        U[        5      (       a¢  SUs=::  a  S:  a  O  O[
        R                  nOrSUs=::  a  S:  a  O  O[
        R                  nOQSUs=::  a  S:  a  O  O[
        R                  nO0SUs=::  a  S:  a  O  O[
        R                  nO[        SU S35      eU R                  XS	9$ [        U[        5      (       ar  S
nSSS-  -  n[        S   " U5      nU[        S5      :X  d  US:X  d  X:w  d  XFs=::  a  U::  a  O  O[
        R                   nO[
        R"                  nU R                  XS	9$ [        U[
        R$                  5      (       a  U R'                  UR(                  5      $ [        XR                  5      (       a  U$ U(       a  [+        SU S[-        U5       S35      eU$ )Né   €ì        l        ìûÿÿÿ         ì            l            zNonrepresentable integer Ú.©rB   g      8g   àÿÿÿ?r   é   ÚabsÚinfg        zcannot convert z	 of type z
 to tensor)Ú
isinstanceÚboolr+   r,   Úget_int1r5   Úint1Úintr6   Úuint32Úint64Úuint64r3   Úscalar_constantÚfloatÚ__builtins__rU   rW   Ú	constexprÚ	to_tensorrR   rE   Útype)r   ÚxÚ
check_typerB   Úmin_float32Úmax_float32Úabs_xs          r   rz   Ú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›Ó$Ø˜‹|Ø‹vØÕ2 {Ö2ÜŸ
™
‘äŸ
™
�Ø×'Ñ'¨Ð'Ð7Ð7ä˜œ2Ÿ<™<×(Ñ(Ø—>‘> !§'¡'Ó*Ð*Ü˜Ÿ;™;×'Ñ'ØˆHÞÜ˜o¨a¨S°	¼$¸q»'¸À*ÐMÓNÐNØˆr!   c                óÞ   • UR                  5       (       aX  U(       d  [        X5      eUR                  5       (       a  X:w  a  [        X5      eUR                  5       (       a  [        X5      eg g r.   )Úis_ptrr   Úis_floating)r   r   r   Úallow_ptr_as       r   Úcheck_ptr_type_implÚ"TritonSemantic.check_ptr_type_implž   sZ   € Ø�=‰=�?‰?ÞÜ/°Ó?Ð?à�}‰}�‰ FÓ$4Ü/°Ó?Ð?à×!Ñ!×#Ñ#Ü/°Ó?Ð?ð $ð r!   c                ó   • [        U[        R                  5      n[        U[        R                  5      nU(       a  Un	U R                  U5      nU(       a  Un
U R                  U5      nUR                  R
                  nUR                  R
                  nU R                  X¼U5        U R                  XËU5        U(       Gay  UR                  5       (       Gdc  UR                  5       (       GdM  U R                  X·XÈU5      nU(       a  W	S:  a  UR                  5       (       d"  U(       a&  W
S:  a   UR                  5       (       a  [        S5      eUR                  5       (       a†  U(       a<  UR                  5       W	s=::  a  UR                  5       ::  d  O  [        SU	 SU 35      eU(       a<  UR                  5       W
s=::  a  UR                  5       ::  d  O  [        SU
 SU 35      eU(       a  U R                  W	US9OU R                  X5      nU(       a  U R                  W
US9OU R                  X-5      nU R!                  X5      u  pX4$ )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 rj   )rn   ÚnumbersÚNumberrz   r{   Úscalarr†   rƒ   rb   Úis_int_unsignedr3   r\   Úget_int_min_valueÚget_int_max_valuerv   ÚcastÚbroadcast_impl_value)r   ÚlhsÚrhsÚallow_lhs_ptrÚallow_rhs_ptrÚarithmetic_checkr_   Úlhs_is_scalarÚrhs_is_scalarÚ
lhs_scalarÚ
rhs_scalarÚ
lhs_sca_tyÚ
rhs_sca_tyÚ
ret_sca_tys                 r   Úbinary_op_type_checking_implÚ+TritonSemantic.binary_op_type_checking_impl©   sÿ  € ô # 3¬¯©Ó7ˆÜ" 3¬¯©Ó7ˆÞØˆJØ—.‘. Ó%ˆCÞØˆJØ—.‘. Ó%ˆCð —X‘X—_‘_ˆ
Ø—X‘X—_‘_ˆ
Ø× Ñ  ¸ÔGØ× Ñ  ¸ÔGß J×$5Ñ$5×$7Ò$7À
×@QÑ@Q×@SÒ@SØ×3Ñ3°JÈzÐjtÓuˆJÞ *¨q£.°Z×5OÑ5O×5QÑ5QÞ$¨°a«¸J×<VÑ<V×<XÑ<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Ó[uˆCÞHU�$×&Ñ& z¸Ð&ÑDÐ[_×[dÑ[dÐehÓ[uˆCð ×,Ñ,¨SÓ6‰ˆØˆxˆr!   c                ó:  • UR                   R                  R                  S:¼  d%  U R                  R                  R
                  (       d  g UR                   R                  nUR                   R                  nXE:X  d   eUR                  5       (       d   eU R                  U[        R                  5      nU R                  U[        R                  5      nU" XS5      nUR                  5       nU R                  U[        R                  5      nUR                  5       nU R                  U[        R                  5      nU R                  U R                  Xg5      U R                  Xh5      5      n	SUR                   SUR                    3n
U R#                  XšS 5        g )Né@   Frr   z! overflow detected for operation )r{   r‹   r@   r,   ÚoptionsÚsanitize_overflowr\   r�   r5   rt   rŽ   rv   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_implÚ/TritonSemantic.binary_op_sanitize_overflow_implÎ   s9  € Ø�8‰8�?‰?×'Ñ'¨2Ó-°T·\±\×5IÑ5I×5[×5[ØØ—X‘X—_‘_ˆ
Ø—X‘X—_‘_ˆ
ØÓ'Ð'Ð'Ø× Ñ ×"Ñ"Ð"Ð"Ø�i‰i˜œRŸX™XÓ&ˆØ�i‰i˜œRŸX™XÓ&ˆÙ˜ %Ó(ˆØ×0Ñ0Ó2ˆ	Ø×(Ñ(¨´B·H±HÓ=ˆ	Ø×0Ñ0Ó2ˆ	Ø×(Ñ(¨´B·H±HÓ=ˆ	Ø�y‰y˜Ÿ™¨Ó8¸$×:LÑ:LÈSÓ:\Ó]ˆØ�J×+Ñ+Ð,Ð,MÈi×N`ÑN`ÐMaÐbˆØ×Ñ˜4 dÕ+r!   c                ó‚  • U R                  XSS5      u  pUR                  R                  nUR                  R                  nUR                  5       (       a   UR                  5       (       a  [	        S5      eUR                  5       (       aC  UR                  5       (       d.  X!p!UR                  R                  nUR                  R                  nUR                  5       (       aî  UR
                  nUR                  R                  5       (       aƒ  UR                  R                  S:  ai  UR                  R                  [        R                  5      R                  U R                  5      nU R                  R                  UR
                  US5      nU R                  U R                  R!                  UR
                  U5      UR                  5      $ UR#                  5       (       aJ  U R                  U R                  R%                  UR
                  UR
                  5      UR                  5      $ UR'                  5       (       am  U(       a  U R)                  XU R*                  5        U R                  U R                  R-                  UR
                  UR
                  5      UR                  5      $ [	        SU 35      e)NTzcannot add pointers togetherr    FrO   )r�   r{   r‹   rƒ   rE   ÚhandlerB   rŒ   r@   Úwith_element_tyr5   rt   Úto_irr,   Úcreate_int_castr+   Úcreate_addptrr„   Úcreate_faddr\   r­   ÚaddÚ
create_add)r   ÚinputÚotherr¢   Úinput_scalar_tyÚother_scalar_tyÚother_handleÚi64_tys           r   r¶   ÚTritonSemantic.addà   sù  € à×8Ñ8¸ÀtÈTÓR‰ˆØŸ*™*×+Ñ+ˆØŸ*™*×+Ñ+ˆØ×!Ñ!×#Ñ#¨×(>Ñ(>×(@Ñ(@ÜÐ:Ó;Ð;ð ×!Ñ!×#Ñ#¨O×,BÑ,B×,DÑ,DØ �5Ø#Ÿ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ÀDÇHÁHÔMØ—;‘;˜tŸ|™|×6Ñ6°u·|±|ÀUÇ\Á\ÓRÐTY×T^ÑT^Ó_Ð_ÜÐ*¨?Ð*;Ð<Ó=Ð=r!   c                óž  • U R                  XSS5      u  pUR                  R                  nUR                  5       (       a  U R	                  XR                  U5      SS9$ UR                  5       (       aJ  U R                  U R                  R                  UR                  UR                  5      UR                  5      $ UR                  5       (       am  U(       a  U R                  XU R                  5        U R                  U R                  R                  UR                  UR                  5      UR                  5      $ [        SU 35      e)NTF)r¢   rO   )r�   r{   r‹   rƒ   r¶   Úminusr„   r+   r,   Úcreate_fsubr°   r\   r­   ÚsubÚ
create_subrE   ©r   r¸   r¹   r¢   r`   s        r   rÂ   ÚTritonSemantic.subÿ   sü   € à×8Ñ8¸ÀtÈUÓS‰ˆØ—J‘J×%Ñ%ˆ	à×Ñ×ÑØ—8‘8˜E§:¡:¨eÓ#4È�8ÐNÐNà× Ñ ×"Ñ"Ø—;‘;˜tŸ|™|×7Ñ7¸¿¹ÀeÇlÁlÓSÐUZ×U_ÑU_Ó`Ð`à×Ñ×ÑÞ Ø×5Ñ5°eÀDÇHÁHÔMØ—;‘;˜tŸ|™|×6Ñ6°u·|±|ÀUÇ\Á\ÓRÐTY×T^ÑT^Ó_Ð_ÜÐ*¨9¨+Ð6Ó7Ð7r!   c                ó2  • U R                  X5      u  pUR                  R                  nUR                  5       (       aJ  U R	                  U R
                  R                  UR                  UR                  5      UR                  5      $ UR                  5       (       am  U(       a  U R                  XU R                  5        U R	                  U R
                  R                  UR                  UR                  5      UR                  5      $ [        SU 35      e©NrO   )r�   r{   r‹   r„   r+   r,   Úcreate_fmulr°   r\   r­   ÚmulÚ
create_mulrE   rÄ   s        r   rÉ   ÚTritonSemantic.mul  sÍ   € à×8Ñ8¸ÓF‰ˆØ—J‘J×%Ñ%ˆ	à× Ñ ×"Ñ"Ø—;‘;˜tŸ|™|×7Ñ7¸¿¹ÀeÇlÁlÓSÐUZ×U_ÑU_Ó`Ð`à×Ñ×ÑÞ Ø×5Ñ5°eÀDÇHÁHÔMØ—;‘;˜tŸ|™|×6Ñ6°u·|±|ÀUÇ\Á\ÓRÐTY×T^ÑT^Ó_Ð_ÜÐ*¨9¨+Ð6Ó7Ð7r!   c                óÐ  • U R                  XSSSS5      u  pUR                  R                  nUR                  R                  nUR                  5       (       a(  UR	                  5       (       a  U R                  X#5      nGOUR	                  5       (       a'  UR                  5       (       a  U R                  X5      nOáUR	                  5       (       aV  UR	                  5       (       aA  U R                  U[        R                  5      nU R                  U[        R                  5      nOvUR                  5       (       aS  UR                  5       (       a>  UR                  UR                  :”  a  U R                  X#5      nO U R                  X5      nO[        SU 35      eU R                  U R                  R                  UR                  UR                  5      UR                  5      $ ©NFTrO   )r�   r{   r‹   r„   r\   r�   r5   rU   Úfp_mantissa_widthrE   r+   r,   Úcreate_fdivr°   )r   r¸   r¹   rº   r»   s        r   ÚtruedivÚTritonSemantic.truediv  sk  € Ø×8Ñ8¸ÀuÈeÐUYÐ[_Ó`‰ˆØŸ*™*×+Ñ+ˆØŸ*™*×+Ñ+ˆà×&Ñ&×(Ñ(¨_×-CÑ-C×-EÑ-EØ—I‘I˜eÓ5ŠEà×#Ñ#×%Ñ%¨/×*EÑ*E×*GÑ*GØ—I‘I˜eÓ5‰Eà×#Ñ#×%Ñ%¨/×*@Ñ*@×*BÑ*BØ—I‘I˜e¤R§Z¡ZÓ0ˆEØ—I‘I˜e¤R§Z¡ZÓ0‰Eà×(Ñ(×*Ñ*¨×/JÑ/J×/LÑ/LØ×0Ñ0°?×3TÑ3TÓTØŸ	™	 %Ó9‘àŸ	™	 %Ó9‘ô Ð.¨Ð.?Ð@ÓAÐAØ�{‰{˜4Ÿ<™<×3Ñ3°E·L±LÀ%Ç,Á,ÓOÐQV×Q[ÑQ[Ó\Ð\r!   c                ó°  • U R                  XSSSS5      u  pUR                  R                  nUR                  R                  nUR                  5       (       añ  UR                  5       (       aÜ  U R	                  X45      nU R                  X5      nU R                  X%5      nUR                  5       (       aJ  U R                  U R                  R                  UR                  UR                  5      UR                  5      $ U R                  U R                  R                  UR                  UR                  5      UR                  5      $ [        SU 35      erÍ   )r�   r{   r‹   r\   rL   r�   Úis_int_signedr+   r,   Úcreate_sdivr°   Úcreate_udivrE   )r   r¸   r¹   rº   r»   Úret_tys         r   ÚfloordivÚTritonSemantic.floordiv7  s
  € Ø×8Ñ8¸ÀuÈeÐUYÐ[_Ó`‰ˆØŸ*™*×+Ñ+ˆØŸ*™*×+Ñ+ˆØ×!Ñ!×#Ñ#¨×(>Ñ(>×(@Ñ(@Ø×.Ñ.¨ÓPˆFØ—I‘I˜eÓ,ˆEØ—I‘I˜eÓ,ˆEØ×#Ñ#×%Ñ%Ø—{‘{ 4§<¡<×#;Ñ#;¸E¿L¹LÈ%Ï,É,Ó#WÐY^×YcÑYcÓdÐdà—{‘{ 4§<¡<×#;Ñ#;¸E¿L¹LÈ%Ï,É,Ó#WÐY^×YcÑYcÓdÐdÜÐ*¨?Ð*;Ð<Ó=Ð=r!   c                óˆ  • UR                   R                  nUR                   R                  nUR                  5       (       a  UR                  5       (       d  [        S5      eU R	                  XSSSS5      u  pU R
                  R                  UR                  UR                  5      nU R                  XaR                   5      $ )Nz4both operands of fdiv must have floating scalar typeFT)	r{   r‹   r„   rE   r�   r,   rÏ   r°   r+   )r   r¸   r¹   Úieee_roundingrº   r»   r¨   s          r   ÚfdivÚTritonSemantic.fdivE  s–   € ØŸ*™*×+Ñ+ˆØŸ*™*×+Ñ+ˆØ×*Ñ*×,Ñ,°O×4OÑ4O×4QÑ4QÜÐRÓSÐSØ×8Ñ8¸ÀuÈeÐUZÐ\`Óa‰ˆØ�l‰l×&Ñ& u§|¡|°U·\±\ÓBˆØ�{‰{˜3§
¡
Ó+Ð+r!   c                óx  • U R                  XSSSS5      u  pUR                  R                  nUR                  R                  nUR                  5       (       aJ  U R	                  U R
                  R                  UR                  UR                  5      UR                  5      $ UR                  5       (       aö  UR                  UR                  :w  a3  [        SUR                  5       -   S-   UR                  5       -   S-   5      eUR                  5       (       aJ  U R	                  U R
                  R                  UR                  UR                  5      UR                  5      $ U R	                  U R
                  R                  UR                  UR                  5      UR                  5      $ [        SU 35      e)NFTzCannot mod z by rP   rO   )r�   r{   r‹   r„   r+   r,   Úcreate_fremr°   r\   rA   rE   r   rÓ   Úcreate_sremÚcreate_urem)r   r¸   r¹   r`   r»   s        r   ÚmodÚTritonSemantic.modN  se  € Ø×8Ñ8¸ÀuÈeÐUYÐ[_Ó`‰ˆØ—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                ó¾  • U R                  X5      u  pUR                  nUR                  5       (       aÞ  U[        R                  R
                  :X  aJ  U R                  U R                  R                  UR                  UR                  5      UR                  5      $ U[        R                  R                  :X  aJ  U R                  U R                  R                  UR                  UR                  5      UR                  5      $ [        SU 35      eUR                  5       (       aJ  U R                  U R                  R                  UR                  UR                  5      UR                  5      $ UR!                  5       (       aJ  U R                  U R                  R#                  UR                  UR                  5      UR                  5      $ [%        SU 35      e©NzUnexpected propagate_nan úUnexpected dtype )r�   rB   r„   r5   ÚPropagateNanÚALLr+   r,   Úcreate_minimumfr°   r{   ÚNONEÚcreate_minnumfr3   rÓ   Úcreate_minsirŒ   Úcreate_minuirE   ©r   r|   ÚyÚpropagate_nanrB   s        r   ÚminimumÚTritonSemantic.minimume  óR  € Ø×0Ñ0°Ó6‰ˆØ—‘ˆØ×Ñ×ÑØ¤§¡× 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                ó¾  • U R                  X5      u  pUR                  nUR                  5       (       aÞ  U[        R                  R
                  :X  aJ  U R                  U R                  R                  UR                  UR                  5      UR                  5      $ U[        R                  R                  :X  aJ  U R                  U R                  R                  UR                  UR                  5      UR                  5      $ [        SU 35      eUR                  5       (       aJ  U R                  U R                  R                  UR                  UR                  5      UR                  5      $ UR!                  5       (       aJ  U R                  U R                  R#                  UR                  UR                  5      UR                  5      $ [%        SU 35      erä   )r�   rB   r„   r5   ræ   rç   r+   r,   Úcreate_maximumfr°   r{   ré   Úcreate_maxnumfr3   rÓ   Úcreate_maxsirŒ   Úcreate_maxuirE   rí   s        r   ÚmaximumÚTritonSemantic.maximumv  rò   r!   c                ó€  • U R                  X#5      u  p#U R                  X5      u  pU R                  X5      u  pUR                  nUR                  5       (       aV  U R                  U R                  R                  UR                  UR                  UR                  U5      UR                  5      $ [        SU S35      e)Nrå   z(. Only floating point clamp is supported)	r�   rB   r„   r+   r,   Úcreate_clampfr°   r{   rE   )r   r|   ÚminÚmaxrï   rB   s         r   ÚclampÚTritonSemantic.clamp‡  s£   € Ø×4Ñ4°SÓ>‰ˆØ×2Ñ2°1Ó:‰ˆØ×2Ñ2°1Ó:‰ˆà—‘ˆØ×Ñ×ÑØ—;‘;˜tŸ|™|×9Ñ9¸!¿(¹(ÀCÇJÁJÐPS×PZÑPZÐ\iÓjÐlm×lrÑlrÓsÐsäÐ/°¨wÐ6^Ð_Ó`Ð`r!   c                ój  • U R                  X5      u  pUR                  R                  nUR                  R                  nUR                  5       (       a  UR                  5       (       d  [	        X45      eU R                  X45      nXS:w  a  U R                  X5      nXT:w  a  U R                  X%5      nX4$ r.   )r�   r{   r‹   r\   r   rL   r�   )r   r¸   r¹   Úinput_sca_tyÚother_sca_tyrœ   s         r   Úbitwise_op_type_checking_implÚ,TritonSemantic.bitwise_op_type_checking_impl–  sš   € Ø×8Ñ8¸ÓF‰ˆØ—z‘z×(Ñ(ˆØ—z‘z×(Ñ(ˆØ×"Ñ"×$Ñ$¨L×,?Ñ,?×,AÑ,AÜ+¨LÓGÐGØ×.Ñ.¨|ÓJˆ
ØÓ%Ø—I‘I˜eÓ0ˆEØÓ%Ø—I‘I˜eÓ0ˆEØˆ|Ðr!   c                ó¼   • U R                  X5      u  pU R                  U R                  R                  UR                  UR                  5      UR
                  5      $ r.   )r  r+   r,   Ú
create_andr°   r{   ©r   r¸   r¹   s      r   r£   ÚTritonSemantic.and_£  óE   € Ø×9Ñ9¸%ÓG‰ˆØ�{‰{˜4Ÿ<™<×2Ñ2°5·<±<ÀÇÁÓNÐPU×PZÑPZÓ[Ð[r!   c                ó¼   • U R                  X5      u  pU R                  U R                  R                  UR                  UR                  5      UR
                  5      $ r.   )r  r+   r,   Ú	create_orr°   r{   r  s      r   Úor_ÚTritonSemantic.or_§  sB   € Ø×9Ñ9¸%ÓG‰ˆØ�{‰{˜4Ÿ<™<×1Ñ1°%·,±,ÀÇÁÓMÈuÏzÉzÓZÐZr!   c                ó¼   • U R                  X5      u  pU R                  U R                  R                  UR                  UR                  5      UR
                  5      $ r.   )r  r+   r,   Ú
create_xorr°   r{   r  s      r   Úxor_ÚTritonSemantic.xor_«  r	  r!   c                ó   • UR                   R                  5       (       d   U R                  U[        R                  5      nUR                   R                  5       (       d   U R                  U[        R                  5      nU R                  X5      $ r.   )r{   Úis_int1Úbitcastr5   rq   r£   r  s      r   Úlogical_andÚTritonSemantic.logical_and¯  s_   € Ø�z‰z×!Ñ!×#Ñ#Ø—L‘L ¬¯©Ó0ˆEØ�z‰z×!Ñ!×#Ñ#Ø—L‘L ¬¯©Ó0ˆEØ�y‰y˜Ó&Ð&r!   c                ó   • UR                   R                  5       (       d   U R                  U[        R                  5      nUR                   R                  5       (       d   U R                  U[        R                  5      nU R                  X5      $ r.   )r{   r  r  r5   rq   r  r  s      r   Ú
logical_orÚTritonSemantic.logical_or¶  s_   € Ø�z‰z×!Ñ!×#Ñ#Ø—L‘L ¬¯©Ó0ˆEØ�z‰z×!Ñ!×#Ñ#Ø—L‘L ¬¯©Ó0ˆEØ�x‰x˜Ó%Ð%r!   c                ó¢   • UR                   R                  5       (       d   U R                  U[        R                  5      nU R                  U5      $ r.   )r{   r  r  r5   rq   Úinvert©r   r¸   s     r   Únot_ÚTritonSemantic.not_½  s8   € Ø�z‰z×!Ñ!×#Ñ#Ø—L‘L ¬¯©Ó0ˆEØ�{‰{˜5Ó!Ð!r!   c                ó¼   • U R                  X5      u  pU R                  U R                  R                  UR                  UR                  5      UR
                  5      $ r.   )r  r+   r,   Úcreate_lshrr°   r{   r  s      r   ÚlshrÚTritonSemantic.lshrÂ  óE   € Ø×9Ñ9¸%ÓG‰ˆØ�{‰{˜4Ÿ<™<×3Ñ3°E·L±LÀ%Ç,Á,ÓOÐQV×Q[ÑQ[Ó\Ð\r!   c                ó¼   • U R                  X5      u  pU R                  U R                  R                  UR                  UR                  5      UR
                  5      $ r.   )r  r+   r,   Úcreate_ashrr°   r{   r  s      r   ÚashrÚTritonSemantic.ashrÆ  r#  r!   c                ó¼   • U R                  X5      u  pU R                  U R                  R                  UR                  UR                  5      UR
                  5      $ r.   )r  r+   r,   Ú
create_shlr°   r{   r  s      r   ÚshlÚTritonSemantic.shlÊ  r	  r!   c                ó   • U$ r.   © r  s     r   ÚplusÚTritonSemantic.plusÒ  s   € Øˆr!   c                óB  • UR                   R                  nUR                  5       (       a  [        SUR	                  5       -   S-   5      eU R                  U R                  R                  UR                  U R                  5      5      U5      nU R                  X1S5      $ )Nz$wrong type argument to unary minus (Ú)T)
r{   r‹   rƒ   r3   r   r+   r,   Úget_null_valuer²   rÂ   )r   r¸   r  Ú_0s       r   rÀ   ÚTritonSemantic.minusÕ  s€   € Ø—z‘z×(Ñ(ˆØ×Ñ× Ñ ÜÐCÀl×F[ÑF[ÓF]Ñ]Ð`cÑcÓdÐdØ�[‰[˜Ÿ™×4Ñ4°\×5GÑ5GÈÏÉÓ5UÓVÐXdÓeˆØ�x‰x˜ 4Ó(Ð(r!   c                ój  • UR                   R                  nUR                  5       (       d  UR                  5       (       a  [	        SUR                  5       -   S-   5      eU R                  U R                  R                  UR                  U R                  5      5      U5      nU R                  X5      $ )Nz%wrong type argument to unary invert (r1  )r{   r‹   rƒ   r„   r3   r   r+   r,   Úget_all_ones_valuer²   r  )r   r¸   r  Ú_1s       r   r  Ú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   • UR                   R                  [        R                  5      $ r.   )r{   r±   r5   rq   )r   Úvs     r   Ú
_bool_likeÚTritonSemantic._bool_likeç  s   € Ø�v‰v×%Ñ%¤b§g¡gÓ.Ð.r!   c                óÈ  • U R                  X5      u  pUR                  R                  nUR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ UR                  5       (       a³  UR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ [        SU 35      erÇ   )r�   r{   r‹   r„   r+   r,   Úcreate_fcmpOGTr°   r;  r\   rÓ   Úcreate_icmpSGTÚcreate_icmpUGTrE   ©r   r¸   r¹   r`   s       r   Úgreater_thanÚTritonSemantic.greater_thanê  ó  € Ø×8Ñ8¸ÓF‰ˆØ—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                óÈ  • U R                  X5      u  pUR                  R                  nUR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ UR                  5       (       a³  UR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ [        SU 35      erÇ   )r�   r{   r‹   r„   r+   r,   Úcreate_fcmpOGEr°   r;  r\   rÓ   Úcreate_icmpSGEÚcreate_icmpUGErE   rA  s       r   r¥   ÚTritonSemantic.greater_equalø  rD  r!   c                óÈ  • U R                  X5      u  pUR                  R                  nUR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ UR                  5       (       a³  UR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ [        SU 35      erÇ   )r�   r{   r‹   r„   r+   r,   Úcreate_fcmpOLTr°   r;  r\   rÓ   Úcreate_icmpSLTÚcreate_icmpULTrE   rA  s       r   Ú	less_thanÚTritonSemantic.less_than  rD  r!   c                óÈ  • U R                  X5      u  pUR                  R                  nUR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ UR                  5       (       a³  UR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ [        SU 35      erÇ   )r�   r{   r‹   r„   r+   r,   Úcreate_fcmpOLEr°   r;  r\   rÓ   Úcreate_icmpSLEÚcreate_icmpULErE   rA  s       r   r¤   ÚTritonSemantic.less_equal  rD  r!   c                ó   • U R                  X5      u  pUR                  R                  nUR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ UR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ [        SU 35      erÇ   )r�   r{   r‹   r„   r+   r,   Úcreate_fcmpOEQr°   r;  r\   Úcreate_icmpEQrE   rA  s       r   ÚequalÚTritonSemantic.equal"  óÂ   € Ø×8Ñ8¸ÓF‰ˆØ—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                ó   • U R                  X5      u  pUR                  R                  nUR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ UR                  5       (       aO  U R	                  U R
                  R                  UR                  UR                  5      U R                  U5      5      $ [        SU 35      erÇ   )r�   r{   r‹   r„   r+   r,   Úcreate_fcmpUNEr°   r;  r\   Úcreate_icmpNErE   rA  s       r   Ú	not_equalÚTritonSemantic.not_equal-  rZ  r!   N)rÖ   c               ó  • [        U[        5      (       a  [        U[        5      (       d  [        S5      e[        US-	  5      n[        US-	  5      nU(       d  U(       a  [        S5      eX!::  a  [        S5      eX!-
  nXfS-
  -  S:w  a  [        S5      eU/nUc%  [        R
                  " [        R                  U5      nUR                  U R                  5      nU R                  U R                  R                  X�U5      U5      $ )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)rn   rr   r3   ro   r5   Ú
block_typer6   r²   r,   r+   Úcreate_make_range)	r   ÚstartÚendrÖ   Úis_start_int64Úis_end_int64ÚrangeÚshapeÚ	ret_ty_irs	            r   ÚarangeÚTritonSemantic.arange<  sÛ   € Ü˜%¤×%Ñ%¬Z¸¼S×-AÑ-AÜÐNÓOÐOÜ˜e r™kÓ*ˆÜ˜C 2™I“ˆÞž\ÜÐ7Ó8Ð8Ø‹<ÜÐ\Ó]Ð]Ø‘ˆØ˜Q‘YÑ AÓ%ÜÐBÓCÐCØ�ˆØ‰>Ü—]’]¤2§8¡8¨UÓ3ˆFØ—L‘L §¡Ó.ˆ	Ø�{‰{˜4Ÿ<™<×9Ñ9¸)ÈCÓPÐRXÓYÐYr!   c                ó  • Uc  [        S5      eUS:X  a5  U R                  R                  UR                  U R                  5      5      nO+[	        U R                  SUR
                   35      nU" U5      nU R                  X5      $ )Nz2dtype must be specified when value is not a tensorr   Úget_)r3   r,   r2  r²   ÚgetattrÚnamer+   )r   rR   rB   Úget_value_fns       r   rv   ÚTritonSemantic.scalar_constantN  so   € à‰=ÜÐQÓRÐRØ�A‹:Ø—L‘L×/Ñ/°·±¸D¿L¹LÓ0IÓJ‰Eä" 4§<¡<°4¸¿
¹
°|Ð1DÓEˆLÙ  Ó'ˆEØ�{‰{˜5Ó(Ð(r!   c                óÆ   • [        U[        R                  5      (       a2  UR                  R                  S:X  d   S5       eU R                  X5      $ U R                  X5      $ )Nr   zonly accepts size-1 tensor)rn   r5   r+   ÚnumelrR   r�   rv   )r   rR   rB   s      r   Úmake_scalarÚTritonSemantic.make_scalarY  sQ   € Ü�eœRŸY™Y×'Ñ'Ø—;‘;×$Ñ$¨Ó)ÐGÐ+GÓGÐ)Ø—9‘9˜UÓ*Ð*à×#Ñ# EÓ1Ð1r!   c                óD   • U R                  U R                  X#5      U5      $ r.   )Úsplatru  )r   ri  rR   rB   s       r   ÚfullÚTritonSemantic.full`  s   € Ø�z‰z˜$×*Ñ*¨5Ó8¸%Ó@Ð@r!   c                óP  • UR                   R                  5       (       a   S5       e[        U5      S:X  a  U$ [        R                  " UR
                  U5      nU R                  U R                  R                  UR                  U R                  5      UR                  5      U5      $ )NzCannot splat a block tensorr   )r{   Úis_blockÚlenr5   rb  rB   r+   r,   Úcreate_splatr²   r°   )r   rR   ri  rÖ   s       r   rx  ÚTritonSemantic.splatg  s|   € Ø—:‘:×&Ñ&×(Ñ(ÐGÐ*GÓGÐ(Üˆu‹:˜‹?ØˆLÜ—’˜uŸ{™{¨EÓ2ˆØ�{‰{˜4Ÿ<™<×4Ñ4°V·\±\À$Ç,Á,Ó5OÐQV×Q]ÑQ]Ó^Ð`fÓgÐgr!   c                ó€   • U R                  U R                  R                  UR                  5      UR                  5      $ r.   )r+   r,   Úcreate_unsplatr°   rB   )r   rR   s     r   ÚunsplatÚTritonSemantic.unsplatn  s*   € Ø�{‰{˜4Ÿ<™<×6Ñ6°u·|±|ÓDÀeÇkÁkÓRÐRr!   c                ó,  • SnU H  nXE-  nM	     UR                   R                  U:w  a  [        S5      e[        R                  " UR                   R
                  U5      nU R                  U R                  R                  UR                  X#5      U5      $ )Nr   z:reshape() cannot change total number of elements in tensor)
r{   rt  r3   r5   rb  r‹   r+   r,   Úcreate_reshaper°   )r   r¸   Ú	dst_shapeÚcan_reorderrt  ÚsrÖ   s          r   ÚreshapeÚTritonSemantic.reshapeq  sw   € ØˆÛˆAØ‰JŠEñ à�:‰:×Ñ˜uÓ$ÜÐYÓZÐZÜ—’˜uŸz™z×0Ñ0°)Ó<ˆØ�{‰{˜4Ÿ<™<×6Ñ6°u·|±|ÀYÓ\Ð^dÓeÐer!   c                ó®  • UR                    Vs/ s H  n[        R                  " U5      PM     nnUR                  US5        UR                  R                  5       (       d  U R                  XS9$ [        R                  " UR                  R                  U5      nU R                  U R                  R                  UR                  U5      U5      $ s  snf )Nr   )ri  )ri  r5   Ú_unwrap_if_constexprÚinsertr{   r|  rx  rb  r‹   r+   r,   Úcreate_expand_dimsr°   )r   r¸   r8   r|   r†  rÖ   s         r   Úexpand_dimsÚTritonSemantic.expand_dimsz  s    € Ø9>¿ºÓEº°A”R×,Ò,¨QÖ/¹ˆ	ÐEØ×Ñ˜˜qÔ!à�z‰z×"Ñ"×$Ñ$Ø—:‘:˜e�:Ð5Ð5ä—’˜uŸz™z×0Ñ0°)Ó<ˆØ�{‰{˜4Ÿ<™<×:Ñ:¸5¿<¹<ÈÓNÐPVÓWÐWùò Fs   � Cc                óf  • U(       d   S5       e[        UR                  5      S:X  d   e[        R                  " UR                  R
                  UR                  S   UR                  S   -   /5      nU R                  U R                  R                  UR                  UR                  5      U5      $ )Nz;current implementation of `cat` always may reorder elementsr   r   )
r}  ri  r5   rb  r{   r‹   r+   r,   Ú
create_catr°   )r   r‘   r’   r‡  Úret_types        r   ÚcatÚ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                óR  • U R                  X5      u  pUR                  / :H  nU(       a$  U R                  US5      nU R                  US5      n[        UR                  S   [        R
                  5      (       a  [        R
                  " S5      nOSnUR                  U/-   n[        R                  " UR                  R                  U5      nU R                  U R                  R                  UR                  UR                  5      U5      nU(       a  U R                  US/SS9nU$ )Nr   éÿÿÿÿr   F©r‡  )r�   ri  r�  rn   r5   ry   rb  r{   r‹   r+   r,   Úcreate_joinr°   r‰  )r   ÚaÚbÚ
was_rank_1ÚtwoÚ	new_shaper“  r¨   s           r   ÚjoinÚTritonSemantic.joinŠ  sã   € Ø×(Ñ(¨Ó.‰ˆð —W‘W ‘]ˆ
ÞØ× Ñ   AÓ&ˆAØ× Ñ   AÓ&ˆAä�a—g‘g˜b‘k¤2§<¡<×0Ñ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                ó�  • [        UR                  5      S:”  d   e[        R                  " UR                  S   5      S:X  d   eUR                  S S n[        R                  " UR
                  R                  U5      nU R                  R                  UR                  5      u  pEU R                  XC5      U R                  XS5      4$ )Nr   r—  r   )r}  ri  r5   rŒ  rb  r{   r‹   r,   Úcreate_splitr°   r+   )r   rš  rž  r“  ÚoutLHSÚoutRHSs         r   ÚsplitÚ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                óà  • [        UR                  5      [        U5      :w  a  [        S5      e[        S U 5       5      [	        [        [        U5      5      5      :w  a  [        SU 35      e[        R                  " UR                  R                  U Vs/ s H  o1R                  U   PM     sn5      nU R                  U R                  R                  UR                  U5      U5      $ s  snf )Nz5permute dims must have the same length as input shapec              3  óN   #   • U  H  n[         R                  " U5      v •  M     g 7fr.   )r5   rŒ  )Ú.0Úds     r   Ú	<genexpr>Ú)TritonSemantic.permute.<locals>.<genexpr>±  s   é € Ð;²d°”"×)Ò)¨!×,Ð,²dùs   ‚#%z?permute dims must be a permutation of 0, 1, ..., n-1, but were )r}  ri  r3   ÚsortedÚlistrh  r5   rb  r{   r‹   r+   r,   Úcreate_transr°   )r   r¸   Údimsrª  r“  s        r   ÚpermuteÚ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ÊTÈ·[±[À´^ÉTÑ4RÓSˆØ�{‰{˜4Ÿ<™<×4Ñ4°U·\±\À4ÓHÈ(ÓSÐSùò 5Ss   ÂC+
c                ó8  • UR                   R                  5       (       d  U R                  X5      $ UR                   R                  5       n[	        U5      [	        U5      :w  a  [        SU SU 35      eX#:X  a  U$ [        U5       H1  u  pEX$   U:w  d  M  US:w  d  M  [        SX$    SU SU SU SU 3
5      e   [        R                  " UR                   R                  U5      nU R                  U R                  R                  UR                  U5      U5      $ )Nz!Cannot broadcast, rank mismatch: ú, r   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension ú: )r{   r|  rx  Úget_block_shapesr}  r3   Ú	enumerater5   rb  r‹   r+   r,   Úcreate_broadcastr°   )r   r¸   ri  Ú	src_shapeÚiÚitemrÖ   s          r   Úbroadcast_impl_shapeÚ#TritonSemantic.broadcast_impl_shape·  s  € Ø�z‰z×"Ñ"×$Ñ$Ø—:‘:˜eÓ+Ð+Ø—J‘J×/Ñ/Ó1ˆ	Üˆy‹>œS ›ZÓ'ÜÐ@ÀÀÈ2ÈeÈWÐUÓVÐVØÓØˆLÜ  Ö+‰GˆAØ‰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           	     óÊ  • UR                   nUR                   nUR                  5       (       a‚  UR                  5       (       dm  UR                  UR                  5      nU R	                  U R
                  R                  UR                  U R
                  5      UR                  5      U5      nX4$ UR                  5       (       d‚  UR                  5       (       am  UR                  UR                  5      nU R	                  U R
                  R                  UR                  U R
                  5      UR                  5      U5      nX4$ UR                  5       (       Ga  UR                  5       (       Gaï  UR                  5       nUR                  5       n[        U5      [        U5      :  a¥  [        [        U5      [        U5      5       H‚  nU R	                  U R
                  R                  UR                  S5      [        R                  " UR                  S/UR                  -   5      5      nUR                   nUR                  5       nM„     O¼[        U5      [        U5      :  a¤  [        [        U5      [        U5      5       H‚  nU R	                  U R
                  R                  UR                  S5      [        R                  " UR                  S/UR                  -   5      5      nUR                   nUR                  5       nM„     [        U5      [        U5      :X  d   e/ n[!        U5       Hs  u  pšXi   nU
S:X  a  UR#                  U5        M"  US:X  d  Xº:X  a  UR#                  U
5        M@  [%        S['        U	5      -   S-   ['        U
5      -   S-   ['        U5      -   5      e   XX:w  aW  [        R                  " UR                  U5      nU R	                  U R
                  R)                  UR                  U5      U5      nXh:w  aW  [        R                  " UR                  U5      nU R	                  U R
                  R)                  UR                  U5      U5      nX4$ )Nr   r   z?Cannot make_shape_compatible: incompatible dimensions at index rµ  r   )r{   r|  r±   r‹   r+   r,   r~  r²   r°   r¶  r}  rh  rŽ  r5   rb  Úvaluesr·  Úappendr3   Ústrr¸  )r   r‘   r’   Úlhs_tyÚrhs_tyÚ	lhs_shapeÚ	rhs_shapeÚ_Ú	ret_shaperº  ÚleftÚrightrÖ   s                r   r�   Ú#TritonSemantic.broadcast_impl_valueÇ  ss  € Ø—‘ˆØ—‘ˆð �?‰?×Ñ V§_¡_×%6Ñ%6Ø×+Ñ+¨F¯M©MÓ:ˆFØ—+‘+˜dŸl™l×7Ñ7¸¿¹ÀTÇ\Á\Ó8RÐTW×T^ÑT^Ó_ÐagÓhˆCðV ˆxˆðS —‘×"Ñ" v§¡×'8Ñ'8Ø×+Ñ+¨F¯M©MÓ:ˆFØ—+‘+˜dŸl™l×7Ñ7¸¿¹ÀTÇ\Á\Ó8RÐTW×T^ÑT^Ó_ÐagÓhˆCðN ˆxˆðK �_‰_×Ò 6§?¡?×#4Ò#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Ö/‘�Ø!™�Ø˜1“9Ø×$Ñ$ UÖ+Ø˜q“j e£mØ×$Ñ$ TÖ*ä$ð &1Ü36°q³6ñ&:Ø<@ñ&AÜCFÀtÃ9ñ&MØOVñ&WÜY\Ð]bÓYcñ&dó eð eñ 0ð Ó%ÜŸš v§}¡}°iÓ@�Ø—k‘k $§,¡,×"?Ñ"?ÀÇ
Á
ÈIÓ"VÐX^Ó_�ØÓ%ÜŸš v§}¡}°iÓ@�Ø—k‘k $§,¡,×"?Ñ"?ÀÇ
Á
ÈIÓ"VÐX^Ó_�àˆxˆr!   c                ó¨   • Uc  g US:X  a  [         R                  R                  $ US:X  a  [         R                  R                  $ [	        SU S35      e)NÚrtneÚrtzzInvalid rounding mode: z0. Supported rounding modes are 'rtne' and 'rtz'.)r   ÚROUNDING_MODEÚRTNEÚRTZr3   )r   Úrounding_modes     r   Ú_str_to_rounding_modeÚ$TritonSemantic._str_to_rounding_modeÿ  sU   € ØÑ ØØ˜FÓ"Ü×#Ñ#×(Ñ(Ð(Ø˜EÓ!Ü×#Ñ#×'Ñ'Ð'ÜÐ2°=°/ÐAqÐrÓsÐsr!   c                óR  • UR                   nUR                  5       (       a  UR                  UR                  5      nX2:X  a  U$ UR                  nUR                  nUR	                  5       (       d  UR	                  5       (       a  U R                  X5      $ UR                  nUR                  nXg:w  a&  [        S[        U5      -   S-   [        U5      -   5      eU R                  U R                  R                  UR                  UR                  U R                  5      5      U5      $ )Nz!Cannot bitcast data-type of size z to data-type of size )r{   r|  r±   r‹   rƒ   r�   Úprimitive_bitwidthr3   rÁ  r+   r,   Úcreate_bitcastr°   r²   )r   r¸   Údst_tyÚsrc_tyÚ
src_sca_tyÚ
dst_sca_tyÚsrc_bitsÚdst_bitss           r   r  ÚTritonSemantic.bitcast  sû   € Ø—‘ˆØ�?‰?×ÑØ×+Ñ+¨F¯M©MÓ:ˆFØÓØˆLØ—]‘]ˆ
Ø—]‘]ˆ
Ø×Ñ×Ñ *×"3Ñ"3×"5Ñ"5Ø—9‘9˜UÓ+Ð+à×0Ñ0ˆØ×0Ñ0ˆØÓÜÐ@Ä3ÀxÃ=ÑPð T2ñ 2Ü47¸³MñBó Cð Cà�{‰{˜4Ÿ<™<×6Ñ6°u·|±|ÀVÇ\Á\ÐRV×R^ÑR^ÓE_Ó`ÐbhÓiÐir!   c                ó(  • UR                   nUR                  nUR                  nXV:X  a  U$ UR                  5       (       a  UR                  U5      nU R	                  U5      nSnUR                  5       (       an  UR                  5       (       aY  UR                  UR                  :  a?  Uc  [        R                  R                  nOJU[        R                  R                  :w  a  SnO)Ub&  [        S[        U5      -   S-   [        U5      -   5      eUR                  5       (       d  UR                  5       (       aL  U R                  R                  R                  S5       c   S5       eU R                  R                  S   " XX0S9$ UR!                  5       (       a  UR                  5       (       d1  UR                  5       (       a  UR!                  5       (       d  U(       aP  U R#                  U R                  R%                  UR&                  UR)                  U R                  5      U5      U5      $ UR+                  5       (       a  UR-                  5       (       a*  UR/                  5       (       aE  UR-                  5       (       d0  U R1                  U R1                  U[2        R4                  5      U5      $ UR                  5       =(       a0    UR                  5       =(       a    UR                  UR                  :„  nU(       aO  U R#                  U R                  R7                  UR&                  UR)                  U R                  5      5      U5      $ UR                  5       =(       a0    UR                  5       =(       a    UR                  UR                  :  n	U	(       aO  U R#                  U R                  R9                  UR&                  UR)                  U R                  5      5      U5      $ UR;                  5       (       GaF  UR;                  5       (       Ga0  UR<                  UR<                  :w  d  UR>                  UR>                  :w  aü  URA                  5       =(       a    URC                  5       (       + n
URC                  5       (       ak  URD                  R)                  U R                  5      nU R#                  U R                  RG                  U5      URD                  5      nU RI                  X5      $ U R#                  U R                  RK                  UR&                  UR)                  U R                  5      U
5      U5      $ URM                  5       (       GaI  UR;                  5       (       Ga3  URC                  5       (       ak  URD                  R)                  U R                  5      nU R#                  U R                  RG                  U5      URD                  5      nU RI                  X5      $ URA                  5       (       aO  U R#                  U R                  RO                  UR&                  UR)                  U R                  5      5      U5      $ U R#                  U R                  RQ                  UR&                  UR)                  U R                  5      5      U5      $ UR;                  5       (       aÝ  URM                  5       (       aÈ  URC                  5       (       d  URA                  5       (       dO  U R#                  U R                  RS                  UR&                  UR)                  U R                  5      5      U5      $ U R#                  U R                  RU                  UR&                  UR)                  U R                  5      5      U5      $ URW                  5       (       aã  UR;                  5       (       aÎ  UR<                  nUS:X  aO  U R#                  U R                  RY                  UR&                  UR)                  U R                  5      5      U5      $ US	:X  ag  U RI                  U R1                  U[2        RZ                  5      U R#                  U R                  R]                  S
5      [2        RZ                  5      5      $ UR;                  5       (       ad  URW                  5       (       aO  U R#                  U R                  R_                  UR&                  UR)                  U R                  5      5      U5      $ URW                  5       (       ad  URW                  5       (       aO  U R#                  U R                  Ra                  UR&                  UR)                  U R                  5      5      U5      $  SU SU 35       e)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 )1r{   r‹   r|  r±   rÒ  r„   rÕ  r   rÎ  rÏ  r3   rÁ  Úis_fp8e4b15r,   Úcodegen_fnsÚgetr[   r+   Úcreate_fp_to_fpr°   r²   rY   rX   rZ   r�   r5   rU   Úcreate_fp_truncÚcreate_fp_extr\   r@   rA   rÓ   Úis_boolrB   r2  r^  r³   Úis_standard_floatingÚcreate_fp_to_siÚcreate_fp_to_uiÚcreate_ui_to_fpÚcreate_si_to_fprƒ   Úcreate_ptr_to_intrt   Ú	get_int64Úcreate_int_to_ptrrÖ  )r   r¸   r×  Úfp_downcast_roundingrØ  rÙ  rÚ  Úuse_custom_roundingÚtruncate_fpÚext_fpÚsign_extendÚtyr3  Úbitwidths                 r   r�   ÚTritonSemantic.cast  sV  € Ø—‘ˆØ—]‘]ˆ
Ø—]‘]ˆ
ØÓ#Øˆ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à—<‘<×+Ñ+Ð,BÒCÀEÐSgÑ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Ò#6Ø×#Ñ# z×'>Ñ'>Ó>À*×B[ÑB[Ð_i×_xÑ_xÓBxØ$×2Ñ2Ó4×Q¸Z×=OÑ=OÓ=QÔ9QˆKØ×!Ñ!×#Ñ#Ø—[‘[×&Ñ& t§|¡|Ó4�Ø—[‘[ §¡×!<Ñ!<¸RÓ!@À%Ç+Á+ÓN�Ø—~‘~ eÓ0Ð0à—{‘{ 4§<¡<×#?Ñ#?ÀÇÁÈfÏlÉlÐ[_×[gÑ[gÓNhÐjuÓ#vØ#)ó+ð +ð ×*Ñ*×,Ò,°×1BÑ1B×1DÒ1DØ×!Ñ!×#Ñ#Ø—[‘[×&Ñ& t§|¡|Ó4�Ø—[‘[ §¡×!<Ñ!<¸RÓ!@À%Ç+Á+ÓN�Ø—~‘~ eÓ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Ñ#DØ×!Ñ!×#Ñ#¨:×+CÑ+C×+EÑ+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Ñ#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Ñ#6Ø—;‘;˜tŸ|™|×=Ñ=¸e¿l¹lÈFÏLÉLÐY]×YeÑYeÓLfÓgÐioÓpÐpð ×Ñ×Ñ :×#4Ñ#4×#6Ñ#6Ø—;‘;˜tŸ|™|×:Ñ:¸5¿<¹<ÈÏÉÐVZ×VbÑVbÓIcÓdÐflÓmÐmà8˜ U G¨4°¨xÐ8Ó8ˆur!   c                ó2  • [         R                  R                  nU(       au  US:X  a  [         R                  R                  nU$ US:X  a  [         R                  R                  nU$ US:X  a  [         R                  R
                  nU$ [        SU S35      eU$ )Nz.caú.cgz.cvúCache modifier ú not supported)r   ÚCACHE_MODIFIERré   ÚCAÚCGÚCVr3   ©r   Úcache_modifierÚcaches      r   Ú_str_to_load_cache_modifierÚ*TritonSemantic._str_to_load_cache_modifier‰  s–   € Ü×!Ñ!×&Ñ&ˆÞØ Ó&Ü×)Ñ)×,Ñ,�ð ˆð   5Ó(Ü×)Ñ)×,Ñ,�ð
 ˆð	   5Ó(Ü×)Ñ)×,Ñ,�ð ˆô ! ?°>Ð2BÀ.Ð!QÓRÐRØˆr!   c                óv  • [         R                  R                  nU(       a—  US:X  a  [         R                  R                  nU$ US:X  a  [         R                  R                  nU$ US:X  a  [         R                  R
                  nU$ US:X  a  [         R                  R                  nU$ [        SU S35      eU$ )Nz.wbrù  z.csz.wtrú  rû  )r   rü  ré   ÚWBrþ  ÚCSÚWTr3   r   s      r   Ú_str_to_store_cache_modifierÚ+TritonSemantic._str_to_store_cache_modifier–  s¸   € Ü×!Ñ!×&Ñ&ˆÞØ Ó&Ü×)Ñ)×,Ñ,�ð ˆð   5Ó(Ü×)Ñ)×,Ñ,�ð ˆð   5Ó(Ü×)Ñ)×,Ñ,�ð
 ˆð	   5Ó(Ü×)Ñ)×,Ñ,�ð ˆô ! ?°>Ð2BÀ.Ð!QÓRÐRØˆr!   c                óî   • [         R                  R                  nU(       aS  US:X  a  [         R                  R                  nU$ US:X  a  [         R                  R                  nU$ [        SU S35      eU$ )NÚ
evict_lastÚevict_firstzEviction policy rû  )r   ÚEVICTION_POLICYÚNORMALÚ
EVICT_LASTÚEVICT_FIRSTr3   )r   Úeviction_policyÚevictions      r   Ú_str_to_eviction_policyÚ&TritonSemantic._str_to_eviction_policy¥  su   € Ü×%Ñ%×,Ñ,ˆÞØ ,Ó.Ü×-Ñ-×8Ñ8�ð
 ˆð	 ! MÓ1Ü×-Ñ-×9Ñ9�ð ˆô !Ð#3°OÐ3DÀNÐ!SÓTÐTØˆr!   c                ó¾   • S nU(       aS  US:X  a  [         R                  R                  nU$ US:X  a  [         R                  R                  nU$ [	        SU S35      eU$ )NÚzeroÚnanzPadding option rû  )r   ÚPADDING_OPTIONÚPAD_ZEROÚPAD_NANr3   )r   Úpadding_optionÚpaddings      r   Ú_str_to_padding_optionÚ%TritonSemantic._str_to_padding_option°  sh   € ØˆÞØ Ó'Ü×+Ñ+×4Ñ4�ð
 ˆð	   5Ó(Ü×+Ñ+×3Ñ3�ð ˆô ! ?°>Ð2BÀ.Ð!QÓRÐRØˆr!   c                óv  • [         R                  R                  nU(       a—  US:X  a  [         R                  R                  nU$ US:X  a  [         R                  R                  nU$ US:X  a  [         R                  R                  nU$ US:X  a  [         R                  R
                  nU$ [        SU S35      eU$ )NÚacquireÚreleaseÚacq_relÚrelaxedúMemory semantic rû  )r   ÚMEM_SEMANTICÚACQUIRE_RELEASEÚACQUIREÚRELEASEÚRELAXEDr3   )r   Ú
sem_optionÚsems      r   Ú_str_to_semÚ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                ó2  • [         R                  R                  nU(       au  US:X  a  [         R                  R                  nU$ US:X  a  [         R                  R                  nU$ US:X  a  [         R                  R                  nU$ [        SU S35      eU$ )NÚgpuÚctaÚsysr%  rû  )r   ÚMEM_SYNC_SCOPEÚGPUÚCTAÚSYSTEMr3   )r   Úscope_optionÚscopes      r   Ú_str_to_scopeÚTritonSemantic._str_to_scopeÊ  s–   € Ü×!Ñ!×%Ñ%ˆÞØ˜uÓ$Ü×)Ñ)×-Ñ-�ð ˆð  Ó&Ü×)Ñ)×-Ñ-�ð
 ˆð	  Ó&Ü×)Ñ)×0Ñ0�ð ˆô !Ð#3°L°>ÀÐ!PÓQÐQØˆr!   c                ó¸  • U(       aÎ  [        US5      (       d  U/nU Vs/ s H0  n[        U[        R                  5      (       a  UR                  OUPM2     nnU H3  n[        U[
        5      (       a  SUs=::  a  [        U5      :  a  M0   e   e   [        U5      S:”  d   e[        U5      [        [        U5      5      :X  d   S5       e[        U5      $ gs  snf )NÚ__iter__r   z'Duplicate dimension in `boundary_check`r-  )	Úhasattrrn   r5   ry   rR   rr   r}  Úsetr­  )r   Úboundary_checkÚblock_shapeÚelemÚdims        r   Ú_canonicalize_boundary_checkÚ+TritonSemantic._canonicalize_boundary_check×  sÌ   € ÞÜ˜>¨:×6Ñ6Ø"0Ð!1�ÙaoÓpÒaoÐY]¬J°t¼R¿\¹\×,JÑ,J˜dŸjšjÐPTÒTÑaoˆNÐpÛ%�Ü! #¤s×+Ñ+°°SÕ0K¼3¸{Ó;KÕ0KÐKÐKÐ0KÐKÐKñ &ä�~Ó&¨Ó*Ð*Ð*Ü�~Ó&¬#¬c°.Ó.AÓ*BÓBÐmÐDmÓmÐBÜ˜.Ó)Ð)Øùò qs    7Cc	           
     óð  • Uc  Ub  [        S5      eUR                  R                  R                  n	U	[        R                  :w  d   S5       eU	R                  5       (       a)  U[        R                  R                  :X  a  [        S5      eUR                  R                  n
U R                  XJR                  5       5      nU R                  U R                  R                  UR                  XEXgU5      U
5      $ )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)r3   r{   Ú
element_tyr5   rq   r\   r   r  r  rC  r¶  r+   r,   Úcreate_tensor_pointer_loadr°   )r   ÚptrÚmaskr¹   r?  r  r  r  Úis_volatileÚelt_tyr×  s              r   Ú_load_block_pointerÚ"TritonSemantic._load_block_pointerã  sÒ   € ð Ñ˜uÑ0ÜÐjÓkÐkà—‘×$Ñ$×/Ñ/ˆØœŸ™Ó ÐXÐ"XÓXÐ Ø�=‰=�?‰?˜w¬"×*;Ñ*;×*CÑ*CÓCÜÐ_Ó`Ð`ð —‘×$Ñ$ˆð ×:Ñ:¸>×KbÑKbÓKdÓeˆð �{‰{Ø�L‰L×3Ñ3°C·J±JÀÐY^ÐjuÓvØóð 	r!   c	           
     óž  • UR                   R                  R                  5       (       d'  [        SUR                   R	                  5        S35      eUc  Ub  [        S5      eU(       d  U(       a  [        S5      eUR                   R                  5       (       db  U(       a*  UR                   R                  5       (       a  [        S5      eU(       a*  UR                   R                  5       (       a  [        S5      eUR                   R                  5       (       a,  Ub  U R                  X5      u  pUb  U R                  X5      u  pUR                   R                  n	U	R                  n
U
[        R                  :H  nU(       aA  [        R                  n
[        R                  " X©R                  5      n	U R                  X5      nUb  U R                  X:5      nUR                   R                  5       (       a  UR                   R                  U
5      nOU
nUc8  U R                  U R                   R#                  UR$                  XgU5      U5      nOVU R                  U R                   R'                  UR$                  UR$                  U(       a  UR$                  OS UXx5      U5      nU(       a   U R                  U[        R                  5      nU$ )NúUnsupported ptr type z in `tl.load`z)`other` cannot be provided without `mask`zä`padding_option` or `boundary_check` argument is not supported for loading a tensor ofpointers or loading a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadúEMask argument cannot be block type if pointer argument is not a blockzFOther argument cannot be block type if pointer argument is not a block)r{   r‹   rƒ   r3   r   r|  r�   rH  r5   rq   Úint8Úpointer_typeÚaddress_spacer�   r±   r+   r,   Úcreate_loadr°   Úcreate_masked_load)r   rJ  rK  r¹   r?  r  r  r  rL  Úptr_tyrM  rç  r×  r¨   s                 r   Ú_load_legacyÚTritonSemantic._load_legacyù  s9  € à�x‰x�‰×%Ñ%×'Ñ'ÜÐ4°S·X±X×5FÑ5FÓ5HÐ4IÈÐWÓXÐXð ‰<˜EÑ-ÜÐHÓIÐIÞ–nÜð Xó Yð Yð
 �x‰x× Ñ ×"Ñ"Þ˜Ÿ	™	×*Ñ*×,Ñ,Ü Ð!hÓiÐiÞ˜Ÿ™×,Ñ,×.Ñ.Ü Ð!iÓjÐjð �8‰8×Ñ×ÑØÑØ ×5Ñ5°cÓ@‘	�ØÑ Ø!×6Ñ6°sÓB‘
�ð —‘—‘ˆØ×"Ñ"ˆð œBŸG™GÑ#ˆÞÜ—W‘WˆFÜ—_’_ V×-AÑ-AÓBˆFØ—)‘)˜CÓ(ˆCð ÑØ—I‘I˜eÓ,ˆ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	           
     óH  • U R                  U5      n	U R                  U5      n
U R                  U5      nUR                  R	                  5       (       a=  UR                  R
                  R                  5       (       a  U R                  XX4X¹X¨5      $ U R                  XX4X¹X¨5      $ r.   )	r  r  r  r{   rƒ   rH  r|  rN  rY  )r   rJ  rK  r¹   r?  r  r  r  rL  r  r  r  s               r   ÚloadÚTritonSemantic.load5  s�   € ð ×0Ñ0°Ó@ˆØ×/Ñ/°Ó@ˆØ×-Ñ-¨nÓ=ˆà�8‰8�?‰?×Ñ §¡×!4Ñ!4×!=Ñ!=×!?Ñ!?à×+Ñ+¨C°uÈgÐ^fÓtÐtð ×$Ñ$ S°ÀwÐW_ÓmÐmr!   c                ó˜  • [        U[        R                  5      (       d   e[        UR                  5      n[        U5      U:X  d   SU S[        U5       35       eU R                  USS9nU R                  R                  UR                  X R                  U5      U R                  U5      5      nU R                  XaR                  5      $ )Nú	expected ú offsets, but got F©Úrequire_i64)rn   r5   Útensor_descriptor_baser}  r@  Ú_convert_to_ir_valuesr,   Úcreate_descriptor_loadr°   r  r  r+   rb  )r   ÚdescÚoffsetsr  r  Úndimr|   s          r   Údescriptor_loadÚTritonSemantic.descriptor_loadC  s³   € ä˜$¤× 9Ñ 9×:Ñ:Ð:Ð:Ü�4×#Ñ#Ó$ˆÜ�7‹|˜tÓ#ÐW y°°Ð6HÌÈWËÈÐ%WÓWÐ#à×,Ñ,¨WÀ%Ð,ÐHˆØ�L‰L×/Ñ/°·±¸W×FfÑFfÐguÓFvØ04×0LÑ0LÈ_Ó0]ó_ˆà�{‰{˜1Ÿo™oÓ.Ð.r!   c                óò   • [        U[        R                  5      (       d   e[        UR                  5      n[        U5      U:X  d   SU S[        U5       35       eUR
                  UR                  :X  d   eg )Nr_  r`  )rn   r5   rc  r}  r@  ri  )r   rf  rR   rg  rh  s        r   Úvalidate_store_likeÚ"TritonSemantic.validate_store_likeN  sl   € Ü˜$¤× 9Ñ 9×:Ñ:Ð:Ð:Ü�4×#Ñ#Ó$ˆÜ�7‹|˜tÓ#ÐW y°°Ð6HÌÈWËÈÐ%WÓWÐ#Ø�{‰{˜d×.Ñ.Ó.Ð.Ñ.r!   c                ó  • U R                  XU5        U R                  X!R                  5      nU R                  USS9nU R	                  U R
                  R                  UR                  UR                  U5      [        R                  5      $ ©NFra  )
rl  r�   rB   rd  r+   r,   Úcreate_descriptor_storer°   r5   Úvoid)r   rf  rR   rg  s       r   Údescriptor_storeÚTritonSemantic.descriptor_storeT  sm   € Ø× Ñ  ¨gÔ6à—	‘	˜%§¡Ó,ˆØ×,Ñ,¨WÀ%Ð,ÐHˆØ�{‰{˜4Ÿ<™<×?Ñ?ÀÇÁÈUÏ\É\Ð[bÓcÔeg×elÑelÓmÐmr!   c                óú  • U R                  XU5        UR                  [        R                  [        R                  [        R
                  [        R                  [        R                  [        R                  1;   d   S5       eU R                  USS9n[        R                  R                  nU R                  U R                  R                  XAR                   UR                   U5      [        R"                  5      $ ©NúUnsupported dtypeFra  )rl  rB   r5   rs   r6   ru   rU   rS   rT   rd  r   ÚDESCRIPTOR_REDUCE_KINDÚADDr+   r,   Úcreate_descriptor_reducer°   rq  ©r   rf  rR   rg  rQ   s        r   Údescriptor_atomic_addÚ$TritonSemantic.descriptor_atomic_add[  s¯   € Ø× Ñ  ¨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Ÿ<™<×@Ñ@ÀÇ{Á{ÐTY×T`ÑT`ÐbiÓjÔln×lsÑlsÓtÐtr!   c                óˆ   • [         R                  R                  5       nUR                  S:H  =(       a    UR                  S:¬  $ )NÚcudaéZ   )r   ÚactiveÚget_current_targetÚbackendÚarch)r   Útargets     r   Ú_has_native_tmaÚTritonSemantic._has_native_tmab  s1   € Ü—‘×1Ñ1Ó3ˆØ—‘ &Ñ(×>¨V¯[©[¸BÑ->Ð?r!   c                óT  • U[         R                  [         R                  [         R                  [         R                  [         R
                  [         R                  1;   d   S5       eU[         R
                  [         R                  1;   a  U R                  5       (       d   S5       eg g )Nrv  z-16-bit float types require native tma support)r5   rs   r6   ru   rt   rS   rT   r…  )r   rB   s     r   Ú$_descriptor_atomic_min_max_supportedÚ3TritonSemantic._descriptor_atomic_min_max_supportedf  sp   € ØœŸ™¤B§H¡H¬b¯i©i¼¿¹Ä2Ç:Á:ÌrÏ{É{Ð[Ó[ÐpÐ]pÓpÐ[Ø”R—Z‘Z¤§¡Ð-Ó-Ø×'Ñ'×)Ñ)ÐZÐ+ZÓZÑ)ð .r!   c                óN  • U R                  XU5        U R                  UR                  5        U R                  USS9n[        R
                  R                  nU R                  U R                  R                  XAR                  UR                  U5      [        R                  5      $ ro  )rl  rˆ  rB   rd  r   rw  ÚMINr+   r,   ry  r°   r5   rq  rz  s        r   Údescriptor_atomic_minÚ$TritonSemantic.descriptor_atomic_mink  ó�   € Ø× Ñ  ¨gÔ6Ø×1Ñ1°$·*±*Ô=Ø×,Ñ,¨WÀ%Ð,ÐHˆÜ×(Ñ(×,Ñ,ˆØ�{‰{˜4Ÿ<™<×@Ñ@ÀÇ{Á{ÐTY×T`ÑT`ÐbiÓjÔln×lsÑlsÓtÐtr!   c                óN  • U R                  XU5        U R                  UR                  5        U R                  USS9n[        R
                  R                  nU R                  U R                  R                  XAR                  UR                  U5      [        R                  5      $ ro  )rl  rˆ  rB   rd  r   rw  ÚMAXr+   r,   ry  r°   r5   rq  rz  s        r   Údescriptor_atomic_maxÚ$TritonSemantic.descriptor_atomic_maxr  rŽ  r!   c                ó¾  • U R                  XU5        UR                  [        R                  [        R                  [        R
                  [        R                  1;   d   S5       eU R                  USS9n[        R                  R                  nU R                  U R                  R                  XAR                  UR                  U5      [        R                  5      $ ru  )rl  rB   r5   rs   r6   ru   rt   rd  r   rw  ÚANDr+   r,   ry  r°   rq  rz  s        r   Údescriptor_atomic_andÚ$TritonSemantic.descriptor_atomic_andy  ó    € Ø× Ñ  ¨gÔ6Ø�z‰zœbŸi™i¬¯©´2·9±9¼b¿h¹hÐGÓGÐ\ÐI\Ó\ÐGØ×,Ñ,¨WÀ%Ð,ÐHˆÜ×(Ñ(×,Ñ,ˆØ�{‰{˜4Ÿ<™<×@Ñ@ÀÇ{Á{ÐTY×T`ÑT`ÐbiÓjÔln×lsÑlsÓtÐtr!   c                ó¾  • U R                  XU5        UR                  [        R                  [        R                  [        R
                  [        R                  1;   d   S5       eU R                  USS9n[        R                  R                  nU R                  U R                  R                  XAR                  UR                  U5      [        R                  5      $ ru  )rl  rB   r5   rs   r6   ru   rt   rd  r   rw  ÚORr+   r,   ry  r°   rq  rz  s        r   Údescriptor_atomic_orÚ#TritonSemantic.descriptor_atomic_or€  s    € Ø× Ñ  ¨gÔ6Ø�z‰zœbŸi™i¬¯©´2·9±9¼b¿h¹hÐGÓGÐ\ÐI\Ó\ÐGØ×,Ñ,¨WÀ%Ð,ÐHˆÜ×(Ñ(×+Ñ+ˆØ�{‰{˜4Ÿ<™<×@Ñ@ÀÇ{Á{ÐTY×T`ÑT`ÐbiÓjÔln×lsÑlsÓtÐtr!   c                ó¾  • U R                  XU5        UR                  [        R                  [        R                  [        R
                  [        R                  1;   d   S5       eU R                  USS9n[        R                  R                  nU R                  U R                  R                  XAR                  UR                  U5      [        R                  5      $ ru  )rl  rB   r5   rs   r6   ru   rt   rd  r   rw  ÚXORr+   r,   ry  r°   rq  rz  s        r   Údescriptor_atomic_xorÚ$TritonSemantic.descriptor_atomic_xor‡  r—  r!   c                ó¼  • [        U[        R                  5      (       d   eUS:X  d   S5       eUS:X  d   S5       e[        UR                  5      S:X  d   SUR                   35       eUR                  S   S:X  d   SUR                   35       e[        UR
                  5      S:X  d   S	UR
                   35       eUR
                  S   S
:¼  d   SUR
                   35       eUR                  nSUR                  -  S
-  nUR                  S   U:¼  d   SU SU SUR                  S    35       e[        R                  " UR                  UR
                  S   UR                  S   /5      nU R                  U4SS9S   nU R                  R                  UR                  UR                  X8R                  U R                  5      5      n	U R                  X˜5      $ )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 ra  zdescriptor gather of ú must have at least ú columns, but got Fra  )rn   r5   rc  r}  r@  ri  rB   rÕ  rb  rd  r,   Úcreate_descriptor_gatherr°   r²   r+   )
r   rf  Ú	x_offsetsÚy_offsetr  r  rB   Úmin_colsr{   r|   s
             r   Údescriptor_gatherÚ 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Ó#Ð#r!   c                ó  • [        U[        R                  5      (       d   e[        UR                  5      S:X  d   SUR                   35       eUR                  S   S:X  d   SUR                   35       e[        UR
                  5      S:X  d   SUR                   35       eUR
                  S   S:¼  d   SUR
                   35       eUR                  nS	UR                  -  S-  nUR                  S   U:¼  d   S
U SU SUR                  S    35       eU R                  U4SS9S   nU R                  R                  UR                  UR                  UR                  U5        U R                  S [        R                  5      $ )Nr   r¢  r   r   r£  r¤  r¥  z6descriptor scatter must have at least 8 rows, but got ra  zdescriptor scatter of r¦  r§  Fra  )rn   r5   rc  r}  r@  ri  ÚshapaerB   rÕ  rd  r,   Úcreate_descriptor_scatterr°   r+   rq  )r   rf  rR   r©  rª  rB   r«  s          r   Údescriptor_scatterÚ!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           	     óâ  • Ub  [        S5      eUR                  R                  R                  5       nUR                  R	                  5       (       d  U R                  X'5      nUR                  R	                  5       (       d   S5       eXrR                  R                  5       :X  d&   SU SUR                  R                  5        S35       eUR                  R                  R                  UR                  R                  :X  d@   SUR                  R                  R                   SUR                  R                   S35       eUR                  R                  R                  nU[        R                  :w  d   S5       eU R                  XG5      nU R                  X(5      nU R                  U R                  R                  UR                  UR                  XEU5      [        R                  5      $ )	NrF  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(rG  )r3   r{   rH  r¶  r|  r¼  r5   rq   rC  r�   r+   r,   Úcreate_tensor_pointer_storer°   rq  )	r   rJ  ÚvalrK  r?  r  r  r@  rM  s	            r   Ú_store_block_pointerÚ#TritonSemantic._store_block_pointer»  sÕ  € ð ÑÜÐjÓkÐkð —h‘h×)Ñ)×:Ñ:Ó<ˆØ�x‰x× Ñ ×"Ñ"Ø×+Ñ+¨CÓ=ˆCØ�x‰x× Ñ ×"Ñ"ÐSÐ$SÓSÐ"ØŸ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˜Ó$ˆð �{‰{Ø�L‰L×4Ñ4°S·Z±ZÀÇÁÈ^ÐdlÓmÔoq×ovÑovóxð 	xr!   c           	     óŒ  • UR                   R                  R                  5       (       d'  [        SUR                   R	                  5        S35      eU(       a  [        S5      eUR                   R                  5       (       d[  UR                   R                  5       (       a  [        S5      eU(       a*  UR                   R                  5       (       a  [        S5      eUR                   R                  5       (       ab  UR                  nUc  U R                  X5      u  pOU R                  XU5      u  pnXqR                  :w  a  [        SUR                   SU 35      eUR                   R                  nUR                  n	U	[        R                  :X  aA  [        R                  n	[        R                  " X˜R                  5      nU R                  X5      nU R                  X)5      nUcO  U R                  U R                   R#                  UR$                  UR$                  XV5      [        R&                  5      $ UR                   R                  R)                  5       (       d  [        S5      eU R                  U R                   R+                  UR$                  UR$                  UR$                  XV5      [        R&                  5      $ )	NrQ  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 blockrR  z(Expected pointer argument to have shape z	 but got ú"Mask must have boolean scalar type)r{   r‹   rƒ   r3   r   r|  ri  Úbroadcast_tensorsrH  r5   rq   rS  rT  rU  r�   r+   r,   Úcreate_storer°   rq  rç  Úcreate_masked_store)
r   rJ  rµ  rK  r?  r  r  Ú	ptr_shaperX  rM  s
             r   Ú_store_legacyÚTritonSemantic._store_legacy×  s"  € à�x‰x�‰×%Ñ%×'Ñ'ÜÐ4°S·X±X×5FÑ5FÓ5HÐ4IÈÐXÓYÐYö Üð Eó Fð Fð
 �x‰x× Ñ ×"Ñ"Ø�x‰x× Ñ ×"Ñ"Ü Ð!iÓjÐjÞ˜Ÿ	™	×*Ñ*×,Ñ,Ü Ð!hÓiÐið �8‰8×Ñ×ÑØŸ	™	ˆIØ‰|Ø×1Ñ1°#Ó;‘��Sà!%×!7Ñ!7¸À$Ó!G‘�˜$ØŸI™IÓ%Ü Ð#KÈCÏIÉIÈ;ÐV_Ð`iÐ_jÐ!kÓlÐlà—‘—‘ˆØ×"Ñ"ˆð ”R—W‘WÓÜ—W‘WˆFÜ—_’_ V×-AÑ-AÓBˆFØ—)‘)˜CÓ(ˆCð �i‰i˜Ó$ˆð ‰<Ø—;‘;˜tŸ|™|×8Ñ8¸¿¹ÀSÇZÁZÐQVÓaÔce×cjÑcjÓkÐkØ�y‰y×Ñ×'Ñ'×)Ñ)ÜÐAÓBÐBØ�{‰{˜4Ÿ<™<×;Ñ;¸C¿J¹JÈÏ
É
ÐTX×T_ÑT_ÐafÓqÜŸ7™7ó$ð 	$r!   c                óÈ  • U R                  U5      nU R                  U5      nUR                  R                  5       (       d)  UR                  R                  R                  5       (       a  [        S5      eUR                  R                  5       (       a<  UR                  R                  R                  5       (       a  U R                  XX4Xx5      $ U R                  XX4Xx5      $ )Nú"Cannot store to a constant pointer)r	  r  r{   Úis_constr‹   r3   rƒ   rH  r|  r¶  r¾  )	r   rJ  rµ  rK  r?  r  r  r  r  s	            r   ÚstoreÚTritonSemantic.store  s®   € ð ×1Ñ1°.ÓAˆØ×/Ñ/°Ó@ˆà�8‰8×Ñ×Ñ #§(¡(§/¡/×":Ñ":×"<Ñ"<ÜÐAÓBÐBà�8‰8�?‰?×Ñ §¡×!4Ñ!4×!=Ñ!=×!?Ñ!?à×,Ñ,¨S°tÈUÓ]Ð]ð ×%Ñ% c°ÀeÓVÐVr!   c           	     óh  • U R                  U5      nU R                  U5      nUR                  R                  R                  nUR
                  S;  a  [        S5      eU R                  U R                  R                  UR                  UR                  UR                  XE5      UR                  5      $ )N)r   ra  r    z9atomic_cas only supports elements with width {16, 32, 64})r-  r9  r{   r‹   rH  rÕ  r3   r+   r,   Úcreate_atomic_casr°   )r   rJ  Úcmprµ  r,  r8  rH  s          r   Ú
atomic_casÚTritonSemantic.atomic_cas  sŒ   € Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ—X‘X—_‘_×/Ñ/ˆ
Ø×(Ñ(°Ó<ÜÐXÓYÐYØ�{‰{˜4Ÿ<™<×9Ñ9¸#¿*¹*ÀcÇjÁjÐRU×R\ÑR\Ð^aÓiÐkn×ksÑksÓtÐtr!   c                óœ  • UR                   R                  R                  5       (       d&  [        SUR                   R	                  5       -   5      eUR                   R                  5       (       d)  UR                   R                  R                  5       (       a  [        S5      eUR                   R                  R                  nU[        R                  L a  US:w  a  [        SU-   S-   5      eU[        R                  L a  US:w  a  [        SU-   S-   5      eU[        R                  [        R                  4;   d  UR                  S:  a  [        SU-   S-   [        U5      -   5      eUR                   R                  5       (       aX  Ub)  U R                  X1R                   R!                  5       5      nUb)  U R                  X!R                   R!                  5       5      nU R#                  X!R                   R                  R                  5      nUc¹  U R$                  R'                  S	5      n[        R(                  nUR                   R                  5       (       a^  UR                   R+                  [        R(                  5      nU R$                  R-                  UR/                  U R$                  5      U5      nU R1                  Xg5      nXU4$ )
Nz)Pointer argument of store instruction is rÁ  r¶   Úatomic_z does not support fp16z does not support bf16r   z does not support T)r{   r‹   rƒ   r3   r   rÂ  rH  r5   rS   rT   Úint16Úuint16rÕ  rÁ  r|  r¼  r¶  r�   r,   rp   rq   r±   r~  r²   r+   )r   rJ  rµ  rK  ÚoprH  Úmask_irÚmask_tys           r   Úatom_red_typechecking_implÚ)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°·x±x×7PÑ7PÓ7RÓS�Ø‰Ø×/Ñ/°·X±X×5NÑ5NÓ5PÓQ�Ø�i‰i˜Ÿ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Ø˜ˆ~Ðr!   c                óâ   • UR                   R                  n[        R                  " USS9nU R	                  X5      nU R                  XBS-
  5      nU R                  U[        R                  5      $ )NF)rö  Úsignedr   )rB   rÕ  r5   Úget_int_dtyper  r!  r�   rq   )r   r|   rö  ÚidtypeÚixÚsignbits         r   Ú_signbitÚTritonSemantic._signbit?  sX   € Ø—7‘7×-Ñ-ˆÜ×!Ò!¨8¸EÑBˆØ�\‰\˜!Ó$ˆØ—)‘)˜B¨1¡Ó-ˆØ�y‰y˜¤"§'¡'Ó*Ð*r!   c                óF  • U R                  XUS5      u  pnU R                  U5      nU R                  U5      nUR                  R                  nUR                  5       (       aó  UR                  5       (       ao  U R                  U R                  R                  [        R                  R                  UR                  UR                  UR                  XE5      UR                  5      $ U R                  U R                  R                  [        R                  R                  UR                  UR                  UR                  XE5      UR                  5      $ U[        R                   [        R"                  1;  a  [%        SU 35      eU[        R                   :X  a  [        R&                  O[        R(                  nU R+                  X'5      nU R+                  U[        R,                  " US5      5      n	U[        R                   :X  a  [        R.                  O[        R0                  n
U R+                  X*5      nU R+                  U[        R,                  " U
S5      5      nU R3                  U5      nU R5                  U5      nU R                  U R                  R                  [        R                  R                  U	R                  UR                  U R7                  X>5      R                  XE5      UR                  5      nU R                  U R                  R                  [        R                  R8                  UR                  UR                  U R7                  X=5      R                  XE5      UR                  5      nU R;                  XïU5      nU R+                  UU5      $ )Nrý   z#atomic_max not supported for dtype r   )rÑ  r-  r9  r{   r‹   r\   rÓ   r+   r,   Úcreate_atomic_rmwr   Ú	ATOMIC_OPr�  r°   ÚUMAXr5   rU   rW   rE   r6   rt   r  rT  rs   ru   rÙ  r  r£   ÚUMINÚwhere©r   rJ  rµ  rK  r,  r8  Úsca_tyÚi_typeÚi_valÚi_ptrÚui_typeÚui_valÚui_ptrÚnegÚposÚpos_retÚneg_retr¨   s                     r   Ú
atomic_maxÚTritonSemantic.atomic_maxF  óŸ  € Ø×8Ñ8¸À4ÈÓO‰ˆ�$Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ—‘—‘ˆà�=‰=�?‰?Ø×#Ñ#×%Ñ%Ø—{‘{Ø—L‘L×2Ñ2´2·<±<×3CÑ3CÀSÇZÁZÐQT×Q[ÑQ[Ð]a×]hÑ]hÐjmÓuØ—H‘Hóð ð —{‘{Ø—L‘L×2Ñ2´2·<±<×3DÑ3DÀcÇjÁjÐRU×R\ÑR\Ð^b×^iÑ^iÐknÓvØ—H‘Hóð ð œ"Ÿ*™*¤b§j¡jÐ1Ó1ÜÐAÀ&ÀÐJÓKÐKà#¤r§z¡zÓ1”—’´r·x±xˆØ—‘˜SÓ)ˆØ—‘˜S¤"§/¢/°&¸!Ó"<Ó=ˆØ%¬¯©Ó3”"—)’)¼¿¹ˆØ—‘˜cÓ+ˆØ—‘˜c¤2§?¢?°7¸AÓ#>Ó?ˆØ�m‰m˜CÓ ˆØ�i‰i˜‹nˆØ—+‘+Ø�L‰L×*Ñ*¬2¯<©<×+;Ñ+;¸U¿\¹\È5Ï<É<Ø+/¯9©9°TÓ+?×+FÑ+FÈóTØUZ×U_ÑU_óaˆð —+‘+Ø�L‰L×*Ñ*¬2¯<©<×+<Ñ+<¸f¿m¹mÈVÏ]É]Ø+/¯9©9°TÓ+?×+FÑ+FÈóTØU[×U`ÑU`óbˆð �j‰j˜ wÓ/ˆØ�|‰|˜C Ó(Ð(r!   c                óF  • U R                  XUS5      u  pnU R                  U5      nU R                  U5      nUR                  R                  nUR                  5       (       aó  UR                  5       (       ao  U R                  U R                  R                  [        R                  R                  UR                  UR                  UR                  XE5      UR                  5      $ U R                  U R                  R                  [        R                  R                  UR                  UR                  UR                  XE5      UR                  5      $ U[        R                   [        R"                  1;  a  [%        SU 35      eU[        R                   :X  a  [        R&                  O[        R(                  nU R+                  X'5      nU R+                  U[        R,                  " US5      5      n	U[        R                   :X  a  [        R.                  O[        R0                  n
U R+                  X*5      nU R+                  U[        R,                  " U
S5      5      nU R3                  U5      nU R5                  U5      nU R                  U R                  R                  [        R                  R                  U	R                  UR                  U R7                  X>5      R                  XE5      UR                  5      nU R                  U R                  R                  [        R                  R8                  UR                  UR                  U R7                  X=5      R                  XE5      UR                  5      nU R;                  XïU5      nU R+                  UU5      $ )Nrü   z#atomic_min not supported for dtype r   )rÑ  r-  r9  r{   r‹   r\   rÓ   r+   r,   rÜ  r   rÝ  r‹  r°   rß  r5   rU   rW   rE   r6   rt   r  rT  rs   ru   rÙ  r  r£   rÞ  rà  rá  s                     r   Ú
atomic_minÚTritonSemantic.atomic_minl  rï  r!   c           
     óÜ  • U R                  XUS5      u  pnU R                  U5      nU R                  U5      nUR                  R                  nUR                  5       (       a  [        R                  R                  O[        R                  R                  nU R                  U R                  R                  XqR                  UR                  UR                  XE5      UR                  5      $ )Nr¶   )rÑ  r-  r9  r{   r‹   r„   r   rÝ  ÚFADDrx  r+   r,   rÜ  r°   )r   rJ  rµ  rK  r,  r8  râ  rÎ  s           r   Ú
atomic_addÚTritonSemantic.atomic_add’  s¯   € Ø×8Ñ8¸À4ÈÓO‰ˆ�$Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ—‘—‘ˆØ"(×"4Ñ"4×"6Ñ"6ŒR�\‰\×Ò¼B¿L¹L×<LÑ<LˆØ�{‰{˜4Ÿ<™<×9Ñ9¸"¿j¹jÈ#Ï*É*ÐVZ×VaÑVaÐcfÓnØŸ8™8ó%ð 	%r!   c           
     óP  • U R                  XUS5      u  pnU R                  U5      nU R                  U5      nU R                  U R                  R                  [        R                  R                  UR                  UR                  UR                  XE5      UR                  5      $ )NÚand)rÑ  r-  r9  r+   r,   rÜ  r   rÝ  r”  r°   r{   ©r   rJ  rµ  rK  r,  r8  s         r   Ú
atomic_andÚTritonSemantic.atomic_and›  óŠ   € Ø×8Ñ8¸À4ÈÓO‰ˆ�$Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ�{‰{Ø�L‰L×*Ñ*¬2¯<©<×+;Ñ+;¸S¿Z¹ZÈÏÉÐUY×U`ÑU`ÐbeÓmÐor×owÑowóyð 	yr!   c           
     óP  • U R                  XUS5      u  pnU R                  U5      nU R                  U5      nU R                  U R                  R                  [        R                  R                  UR                  UR                  UR                  XE5      UR                  5      $ )NÚor)rÑ  r-  r9  r+   r,   rÜ  r   rÝ  r™  r°   r{   rù  s         r   Ú	atomic_orÚTritonSemantic.atomic_or¢  sˆ   € Ø×8Ñ8¸À4ÈÓN‰ˆ�$Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ�{‰{Ø�L‰L×*Ñ*¬2¯<©<¯?©?¸C¿J¹JÈÏ
É
ÐTX×T_ÑT_ÐadÓlÐnq×nvÑnvóxð 	xr!   c           
     óP  • U R                  XUS5      u  pnU R                  U5      nU R                  U5      nU R                  U R                  R                  [        R                  R                  UR                  UR                  UR                  XE5      UR                  5      $ )NÚxor)rÑ  r-  r9  r+   r,   rÜ  r   rÝ  r�  r°   r{   rù  s         r   Ú
atomic_xorÚTritonSemantic.atomic_xor©  rü  r!   c           
     óP  • U R                  XUS5      u  pnU R                  U5      nU R                  U5      nU R                  U R                  R                  [        R                  R                  UR                  UR                  UR                  XE5      UR                  5      $ )NÚxchg)rÑ  r-  r9  r+   r,   rÜ  r   rÝ  ÚXCHGr°   r{   rù  s         r   Úatomic_xchgÚTritonSemantic.atomic_xchg°  s†   € Ø×8Ñ8¸À4ÈÓP‰ˆ�$Ø×Ñ˜sÓ#ˆØ×"Ñ" 5Ó)ˆØ�{‰{Ø�L‰L×*Ñ*¬2¯<©<×+<Ñ+<¸c¿j¹jÈ#Ï*É*ÐVZ×VaÑVaÐcfÓnØ�H‰Hóð 	r!   c                ó@  • UR                  5       U R                  R                  R                  ;   d+   SU R                  R                  R                   SU 35       eUR	                  5       nUS:X  a  SnUS:X  a  SnUS:X  a  Sn[        [        R                  U5      $ )	Nzinput_precision must be one of ú. Got ÚTF32X3ÚTF32x3ÚBF16X3ÚBF16x3ÚBF16X6ÚBF16x6)Úlowerr,   r¡   Úallowed_dot_input_precisionsÚupperro  r   ÚINPUT_PRECISION)r   Úinput_precisions     r   Ú_str_to_dot_input_precisionÚ*TritonSemantic._str_to_dot_input_precision¼  sž   € Ø×$Ñ$Ó&¨$¯,©,×*>Ñ*>×*[Ñ*[Ó[ð 	yØ-¨d¯l©l×.BÑ.B×._Ñ._Ð-`Ð`fÐgvÐfwÐxó	yÐ[à)×/Ñ/Ó1ˆØ˜hÓ&Ø&ˆOØ˜hÓ&Ø&ˆOØ˜hÓ&Ø&ˆOÜ”r×)Ñ)¨?Ó;Ð;r!   c           
     óô  • UR                   R                  5       (       a  UR                   R                  5       (       d   eUR                  R                  5       (       a!  UR                  R                  5       (       a  GO7UR                  [        R
                  [        R                  [        R                  [        R                  [        R                  [        R                  4;   d   SUR                   35       eUR                  [        R
                  [        R                  [        R                  [        R                  [        R                  [        R                  4;   d   SUR                   35       eUR                  UR                  :X  d!   SUR                   SUR                   35       eUR                  R                  5       (       d  UR                  R                  5       (       az  SU R                  R                  R                  ;   a  [        R                   " S5        U R#                  U[        R                  5      nU R#                  U[        R                  5      nUR                  R%                  5       =(       d    UR                  R%                  5       nUR                  R'                  5       =(       d    UR                  R'                  5       nU(       d  U(       a­  U(       a  SOSn	X�R                  R                  R                  ;   a  U R                  R                  R(                  n
[        R                   " U	 S	U
 S
U
 S35        U R#                  U[        R                  5      nU R#                  U[        R                  5      nUc   U R                  R                  R*                  nU R-                  U5      n[/        UR0                  5      n[/        UR0                  5      nX¼s=:X  a  S:X  d2  O  X¼s=:X  a  S:X  d$  O   SUR0                   SUR0                   S35       eUR0                  S   R2                  UR0                  S   R2                  :X  dV   SUR0                   SUR0                   SUR0                  S   R2                   SUR0                  S   R2                   S3	5       eU R                  R4                  R7                  S5       c   S5       eU R                  R4                  S   " UR                   UR                   5      nUR0                  S   R2                  US   :¼  a@  UR0                  S   R2                  US   :¼  a   UR0                  S   R2                  US   :¼  d   SUS    SUS    SUS    35       eUR                   R8                  R;                  5       (       a\  UR                   R8                  [        R
                  :X  d   S5       eU R                  R=                  S5      n[        R>                  nGO@URA                  5       (       a  [C        S5      eUR                   R8                  RE                  5       (       d)  UR                   R8                  RA                  5       (       a,  U R                  RG                  S5      n[        R                  nO¢UR                   R8                  RI                  5       (       a,  U R                  RK                  S5      n[        R                  nOMURM                  5       (       a  U R                  RO                  S5      OU R                  RG                  S5      nUnUR                   R0                  S   nUR                   R0                  S   nUR                   R0                  S   nUS:X  a  UR                   R0                  S   OS n[        RP                  " UU(       a  UUU/OUU/5      nUc6  U R                  RS                  URU                  U R                  5      U5      nOLURV                  nUR                   R0                  UR0                  :X  a  UR                   RX                  U:X  d   eUcb  UR                  R                  5       (       a@  UR                  R                  5       (       a!  U R                  R                  RZ                  nOYSnOVUR                  R                  5       (       a7  UR                  R                  5       (       a  UU:”  a  [C        S U S!U S35      eU R]                  U R                  R_                  URV                  URV                  UXE5      U5      $ )"NzUnsupported lhs dtype zUnsupported rhs dtype z&Both operands must be same dtype. Got r   Úfp8e4b15z–the use of fp8e4b15 is deprecated on Hopper and later architectures and can cause significant slow down. It will be removed in a future triton releaseÚfp8e4b8Úfp8e5b16z- is AMD gfx942 specific and not supported on z^ so it's upcasted to fp16 and can cause significant slow down. Please use OCP fp8 variants on z for performancer   é   ú+Both inputs must be either 2D or 3D; (lhs: ú	 vs rhs: r1  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 ()0r{   r|  rB   r[   r5   rS  Úuint8rS   rT   rU   rW   rá  r,   r¡   Ú!deprecated_fp8_dot_operand_dtypesÚwarningsÚwarnr�   Ú
is_fp8e4b8Úis_fp8e5b16rƒ  Údefault_dot_input_precisionr  r}  ri  rR   râ  rã  r‹   r\   Ú	get_int32r6   rZ   r3   rX   Úget_fp32rV   Úget_fp64rY   Úget_fp16rb  r~  r²   r°   rH  Úmax_num_imprecise_acc_defaultr+   Ú
create_dot)r   r‘   r’   Úaccr  Úmax_num_imprecise_accÚ	out_dtypeÚuses_fp8e4b8Úuses_fp8e5b16Ú	type_namerƒ  Úlhs_rankÚrhs_rankr!  r3  Úret_scalar_tyÚMÚNÚKÚBrÖ   Ú
acc_handles                         r   ÚdotÚTritonSemantic.dotÈ  s½  € à�x‰x× Ñ ×"Ñ" s§x¡x×'8Ñ'8×':Ñ':Ð:Ð:à�9‰9×Ñ×Ñ #§)¡)×"2Ñ"2×"4Ñ"4áà—9‘9¤§¡¬"¯(©(´B·J±JÄÇÁÌRÏZÉZÜ!#§¡ð!-ó -ð SØ0FÀsÇyÁyÀkÐ.RóSð -à—9‘9¤§¡¬"¯(©(´B·J±JÄÇÁÌRÏZÉZÜ!#§¡ð!-ó -ð SØ0FÀsÇyÁyÀkÐ.RóSð -à—9‘9 §	¡	Ó)ÐoÐ-SÐTW×T]ÑT]ÐS^Ð^cÐdg×dmÑdmÐcnÐ+oÓoÐ)à�9‰9× Ñ ×"Ñ" c§i¡i×&;Ñ&;×&=Ñ&=Ø˜TŸ\™\×1Ñ1×SÑSÓSÜ—’ð môð —)‘)˜C¤§¡Ó,ˆCØ—)‘)˜C¤§¡Ó,ˆCà—y‘y×+Ñ+Ó-×G°·±×1EÑ1EÓ1GˆØŸ	™	×-Ñ-Ó/×J°3·9±9×3HÑ3HÓ3JˆÞž=Þ%1™	°zˆIØŸL™L×0Ñ0×RÑRÓRØ—|‘|×+Ñ+×0Ñ0�Ü—’Ø �kÐ!NÈtÈfð U6Ø6:°VÐ;KðMôNð —i‘i ¤R§Z¡ZÓ0�Ø—i‘i ¤R§Z¡ZÓ0�àÑ"Ø"Ÿl™l×2Ñ2×NÑNˆOà×:Ñ:¸?ÓKˆä�s—y‘y“>ˆÜ�s—y‘y“>ˆØÕ( 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ð à�|‰|×'Ñ'×+Ñ+ØóØ#'ñ(ð 	^Ø)]ó	^ð (à—|‘|×/Ñ/°Ò?ÀÇÁÈ#Ï(É(Ó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Ñ*CØ—‘×&Ñ& qÓ)ˆBÜŸJ™J‰MØ�X‰X�_‰_×$Ñ$×&Ñ&Ø—‘×&Ñ& qÓ)ˆBÜŸJ™J‰Mà-6×->Ñ->×-@Ñ-@�—‘×&Ñ& qÔ)ÀdÇlÁl×F[ÑF[Ð\]ÓF^ˆBØ%ˆMà�H‰H�N‰N˜2ÑˆØ�H‰H�N‰N˜2ÑˆØ�H‰H�N‰N˜2ÑˆØ!)¨Q£ˆC�H‰H�N‰N˜1Ò°DˆÜ—’˜}¾1¨q°!°Q©iÀ1ÀaÀ&ÓIˆØ‰;ØŸ™×2Ñ2°6·<±<ÀÇÁÓ3MÈrÓR‰JàŸ™ˆJØ—8‘8—>‘> V§\¡\Ó1°c·h±h×6IÑ6IÈYÓ6VÐVÐVð !Ñ(Ø�y‰y×Ñ×!Ñ! c§i¡i×&6Ñ&6×&8Ñ&8Ø(,¯©×(<Ñ(<×(ZÑ(ZÑ%à()Ñ%à�y‰y×Ñ×!Ñ! c§i¡i×&6Ñ&6×&8Ñ&8Ð=RÐUVÓ=VÜ Ð#:Ð;PÐ:QÐQaÐbcÐadÐdeÐ!fÓgÐgà�{‰{Ø�L‰L×#Ñ# C§J¡J°·
±
¸JÈÓoÐqwóyð 	yr!   c                ó|   • [        [        R                  UR                  5       S 5      nUc  [	        SU S35      eU$ )NzInvalid float format: ri   )ro  r   ÚScaleDotElemTypeTYr  r3   )r   Úfloat_formatÚty_enums      r   Ú_str_to_fp_typeÚTritonSemantic._str_to_fp_type%  s>   € Üœ"×/Ñ/°×1CÑ1CÓ1EÀtÓLˆØ‰?ÜÐ5°l°^À1ÐEÓFÐFØˆr!   c                óD  • [         R                  [         R                  [         R                  [         R                  S.R                  U5      nUcD  US:X  d
   SU 35       eUR                  [         R                  :X  d   SUR                   35       eU$ UR                  U:X  a  U$ [         R                  [         R                  [         R                  [         R                  S.U   nUR                  U:X  d   SU SUR                   35       eU R                  X5      $ )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  )
r5   Úfloat8e5Ú
float8e4nvrT   rS   rã  rB   r"  rÍ  r  )r   rµ  rA  Ú	triton_tyÚunsigned_tys        r   Ú_bitcast_to_fp_typeÚ"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Ð+Ø—<‘< Ó/Ð/r!   c                óf  • UbU  UR                   R                  5       (       a  SOSnUR                  R                  nUXU-  /:X  d   SU SX6-   SU 35       eUbV  UR                   R                  5       (       a  SOSnUR                  R                  nUX#U-  /:X  d   SU SX6-   SU 35       eg g )Nr   ra  z%lhs_scale must be a tensor of shape [r´  z]. Got z%rhs_scale must be a tensor of shape [)rB   Ú
is_fp8e4nvr{   ri  )	r   r8  r9  r:  Ú	lhs_scaleÚ	rhs_scaleÚscale_factorÚlhs_scale_shapeÚrhs_scale_shapes	            r   Úverify_scaled_shapeÚ"TritonSemantic.verify_scaled_shape=  sé   € ØÑ Ø!*§¡×!;Ñ!;×!=Ñ!=™2À2ˆLØ'Ÿn™n×2Ñ2ˆOØ"Ø˜Ñ$ð'ó ð hà6°q°c¸¸AÑ<MÐ;NÈgÐVeÐUfÐgóhð ð Ñ Ø!*§¡×!;Ñ!;×!=Ñ!=™2À2ˆLØ'Ÿn™n×2Ñ2ˆOØ"Ø˜Ñ$ð'ó ð hà6°q°c¸¸AÑ<MÐ;NÈgÐVeÐUfÐgóhñ ð !r!   c                óâ  • UR                   R                  5       (       a  UR                   R                  5       (       d   e[        UR                  5      n[        UR                  5      nXÍs=:X  a  S:X  d2  O  XÍs=:X  a  S:X  d$  O   SUR                   SUR                   S35       eUR                  nUR                  nU R                  U5      nU R                  U5      n1 SknUU;   d
   SU 35       eUU;   d
   SU 35       eUS L =(       d/    [        U[        R                  5      =(       a    UR                  S L nUS L =(       d/    [        U[        R                  5      =(       a    UR                  S L nU R                  X5      nU R                  XF5      nU	(       d  US	:X  d   S
5       eU
(       d  US	:X  d   S
5       eUR                   R                  SS  u  nnUR                   R                  SS  u  nnUS	:X  a  SOSnUS	:X  a  SOSnU	(       a  UU-  OUnU
(       a  UU-  OUnUU:X  d"   SUR                   SUR                   S35       eUS:X  a  UR                   R                  S   OS nUnU	(       d  UU-  nOUU-  nU
(       d  UU-  n[        R                  " UU(       a  UUU/OUU/5      nU R                  R                  S5      nUc6  U R                  R                  UR                  U R                  5      U5      nOLUR                  nUR                   R                  UR                  :X  a  UR                   R                   U:X  d   eU(       a  S OUR                  n U(       a  S OUR                  n!U R#                  UUUU(       a  S OUU(       a  S OU5        U R%                  U R                  R'                  UR                  U!XäR                  U XøXšU5
      U5      $ )Nr   r  r  r  r1  >   rH  rJ  rG  rF  rI  zNYI: lhs_format zNYI: rhs_format rJ  zBonly mxfp4 inputs can be packed along a dimension different than Kr   r   zCReduction dimension should pack the same number of elements; (lhs: r   )r{   r|  r}  ri  rR   rC  rn   r5   ry   rO  rb  r,   r*  r~  r²   r°   rH  rX  r+   Úcreate_dot_scaled)"r   r‘   rS  Ú
lhs_formatr’   rT  Ú
rhs_formatr/  Ú	fast_mathÚ
lhs_k_packÚ
rhs_k_packr1  r5  r6  Úlhs_format_enumÚrhs_format_enumÚallowed_formatsÚrhs_scale_is_noneÚlhs_scale_is_noner8  ÚK_LHSÚK_RHSr9  ÚPACKED_AÚPACKED_BÚPACKED_A_DIMÚPACKED_B_DIMr;  r:  rÖ   r3  r<  Úrhs_scale_handleÚlhs_scale_handles"                                     r   Ú
dot_scaledÚTritonSemantic.dot_scaledK  sØ  € ð �x‰x× Ñ ×"Ñ" s§x¡x×'8Ñ'8×':Ñ':Ð:Ð:ä�s—y‘y“>ˆÜ�s—y‘y“>ˆØÕ( 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Ó7ˆØ×&Ñ& sÓ7ˆæ˜Z¨6Ó1ÐwÐ3wÓwÐ1Þ˜Z¨6Ó1ÐwÐ3wÓwÐ1Ø—8‘8—>‘> " #Ð&‰ˆˆ5Ø—8‘8—>‘> " #Ð&‰ˆˆqØ" fÓ,‘1°!ˆØ" fÓ,‘1°!ˆÞ+5�x %Ò'¸5ˆÞ+5�x %Ò'¸5ˆØ˜|Ó+ð  	TÐ/rÐsv×s|Ñs|Ðr}ð  ~Gð  HK÷  HQñ  HQð  GRð  RSð  .Tó  	TÐ+à!)¨Q£ˆC�H‰H�N‰N˜1Ò°DˆØˆÞØ�H‘‰Aà�H‘ˆAÞØ�H‘ˆAÜ—’˜y¶q¨1¨a°©)¸qÀ!¸fÓEˆØ�\‰\×"Ñ" 1Ó%ˆØ‰;ØŸ™×2Ñ2°6·<±<ÀÇÁÓ3MÈrÓR‰JàŸ™ˆJØ—8‘8—>‘> V§\¡\Ó1°c·h±h×6IÑ6IÈYÓ6VÐVÐVÞ#4™4¸)×:JÑ:JÐÞ#4™4¸)×:JÑ:JÐØ× Ñ   A qÖ2C©$ÈÞ):¡À	ô	Kà�{‰{Ø�L‰L×*Ñ*¨3¯:©:Ð7GÈ×ZdÑZdÐfvØ+:ÀzÐ_iókØlrótð 	tr!   c                óN  • UR                   [        R                  :w  a#  [        R                  " SUR                    35        U R                  U[        R                  5      nU R                  X#SS5      u  p#UR                  R                  5       (       a'  U R                  X5      u  pU R                  X#5      u  p#OU R                  X5      u  pUR                  nU R                  U R                  R                  UR                  UR                  UR                  5      U5      $ )Nzgtl.where with a non-boolean condition is deprecated and will error out in a future triton release. Got T)rB   r5   rq   r$  r%  r�   r�   r{   r|  r�   r+   r,   Úcreate_selectr°   )r   Ú	conditionr|   rî   rÆ  rÖ   s         r   rà  ÚTritonSemantic.where„  sù   € Ø�?‰?œbŸg™gÓ%Ü�MŠMØyð  {D÷  {Jñ  {Jð  zKð  Lôð —I‘I˜i¬¯©Ó1ˆ	Ø×0Ñ0°°t¸TÓB‰ˆà�>‰>×"Ñ"×$Ñ$Ø×4Ñ4°YÓB‰LˆIØ×,Ñ,¨QÓ2‰DˆAˆqà×4Ñ4°YÓB‰LˆIØ—‘ˆØ�{‰{˜4Ÿ<™<×5Ñ5°i×6FÑ6FÈÏÉÐRS×RZÑRZÓ[Ð]cÓdÐdr!   c                ód   • U(       a  [         R                  " X#5      nOUnU R                  X5      $ r.   )r5   rb  r+   )r   r|   r`   rÇ  Úres_tys        r   Úwrap_tensorÚTritonSemantic.wrap_tensor˜  s)   € ÞÜ—]’] 9Ó8‰Fð ˆFØ�{‰{˜1Ó%Ð%r!   c                óL  ^ ^^^	^
• Uc  [        U 4S jT 5       5      mSnTS   R                  R                  m
[        T
5      nX$:  d   SU S35       e[	        T
5       VVs/ s H  u  pVXR:w  d  M  UPM     snnm	[        U
4S jT 5       5      (       d   S5       eT R                  R                  T Vs/ s H  owR                  PM     snU5      mU" T5        TR                  5       (       d   e[        UUU	U 4S j[        [        T5      5       5       5      $ s  snnf s  snf )Nc              3  ón   >#   • U  H*  nTR                  XR                  R                  /S S9v •  M,     g7f)Tr˜  N)r‰  rt  rR   )r©  Útr   s     €r   r«  Ú+TritonSemantic.reduction.<locals>.<genexpr>¢  s+   øé € Ð^ÒW]ÐRS˜4Ÿ<™<¨¯G©G¯M©M¨?È˜<ÕMÒW]ùs   ƒ25r   z&reduction axis must be < inputs rank (r1  c              3  óT   >#   • U  H  oR                   R                  T:H  v •  M     g 7fr.   )r{   ri  )r©  rz  ri  s     €r   r«  r{  ©  s   øé € Ð9²&¨Q—6‘6—<‘< 5Ö(²&ùs   ƒ%(z-all reduction inputs must have the same shapec              3  ó–   >#   • U  H>  nTR                  TR                  U5      TU   R                  R                  T5      v •  M@     g 7fr.   ©rv  Ú
get_resultr{   r‹   )r©  rº  ÚinputsÚ	reduce_oprÇ  r   s     €€€€r   r«  r{  ¯  sB   øé € ð uÚasÐ\]ˆD×Ñ˜Y×1Ñ1°!Ó4°f¸Q±i·n±n×6KÑ6KÈY×WÐWÒasùó   ƒAA	)Útupler{   ri  r}  r·  Úallr,   Úcreate_reducer°   Úverifyrh  )r   r€  r8   Úregion_builder_fnÚrankrº  rˆ  rz  r�  rÇ  ri  s   ``      @@@r   Ú	reductionÚTritonSemantic.reduction   s  ü€ Ø‰<ÜÔ^ÑW]Ó^Ó^ˆFØˆDà�q‘	—‘×$Ñ$ˆÜ�5‹zˆØ‹{ÐLÐDÀTÀFÈ!ÐLÓLˆ{Ü#,¨UÔ#3ÔAÒ#3™4˜1°q±y—QÑ#3ÒAˆ	ÜÔ9±&Ó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Á1DÂ5D!c                óþ  ^ ^^^• TS   R                   R                  m[        T5      nU* Us=::  a  U:  d  O   SU SU S35       eUS:  a  X%-  nT H$  nUR                   R                  T:X  a  M   S5       e   T R                  R	                  T Vs/ s H  ofR
                  PM     snX$5      mU" T5        TR                  5       (       d   e[        UUU U4S j[        [        T5      5       5       5      $ s  snf )Nr   z
scan axis z must be < inputs rank (r1  z(all scan inputs must have the same shapec              3  ó–   >#   • U  H>  nTR                  TR                  U5      TU   R                  R                  T5      v •  M@     g 7fr.   r~  )r©  rº  r€  Úscan_opr   ri  s     €€€€r   r«  Ú2TritonSemantic.associative_scan.<locals>.<genexpr>Ç  s@   øé € ÐwÒdvÐ_`�T×%Ñ% g×&8Ñ&8¸Ó&;¸VÀA¹Y¿^¹^×=RÑ=RÐTY×ZÐZÒdvùr‚  )	r{   ri  r}  r,   Úcreate_scanr°   r†  rƒ  rh  )	r   r€  r8   r‡  Úreverserˆ  rz  r�  ri  s	   ``     @@r   Úassociative_scanÚTritonSemantic.associative_scan¶  sç   û€ à�q‘	—‘×$Ñ$ˆÜ�5‹zˆàˆu˜Õ#˜tÕ#ÐW z°$°Ð7OÐPTÈvÐUVÐ%WÓWÐ#à�!‹8Ø‰LˆDãˆAØ—6‘6—<‘< 5Õ(ÐTÐ*TÓTÐ(ñ ð —,‘,×*Ñ*¹fÓ+Eºf¸¯H¬H¹fÑ+EÀtÓUˆÙ˜'Ô"Ø�~‰~×ÑÐÐä×wÔdiÔjmÐntÓjuÔdvÓwÓwÐwùò	 ,Fs   ÂC:c                ó¶  • UR                   R                  5       (       d   S5       e[        UR                  R                  5      n[        UR                  R                  5      U:X  d   S5       eU* Us=::  a  U:  d  O   SU SU S35       eUS:  a  X4-  n[        U5       HI  nXS:X  a  M
  UR                  R                  U   UR                  R                  U   :X  a  M@   SU S35       e   U R                  R                  UR                  UR                  U5      nU R                  XaR                  R                  UR                  R                  5      $ )	Nzindex must be an integer tensorz0source and index tensors must have the same rankzgather axis z must be < source rank (r1  r   z
index dim z( must match the corresponding source dim)rB   r\   r}  r{   ri  rh  r,   Úcreate_gatherr°   rv  r‹   )r   ÚsrcÚindexr8   rˆ  rª  Úgathers          r   r—  ÚTritonSemantic.gatherÍ  s'  € Ø�{‰{×!Ñ!×#Ñ#ÐFÐ%FÓFÐ#ä�3—8‘8—>‘>Ó"ˆÜ�5—:‘:×#Ñ#Ó$¨Ó,Ð`Ð.`Ó`Ð,àˆu˜Õ#˜tÕ#ÐY |°D°6Ð9QÐRVÐQWÐWXÐ%YÓYÐ#Ø�!‹8Ø‰LˆDä�t–ˆAØ‹yÙØ—:‘:×#Ñ# AÑ&¨#¯(©(¯.©.¸Ñ*;Õ;Ðx¸zÈ$ÈÐOwÐ=xÓxÐ;ñ ð
 —‘×+Ñ+¨C¯J©J¸¿¹ÀdÓKˆØ×Ñ ¯©¯©¸¿¹×9IÑ9IÓJÐJr!   c                óú   • U(       d  gUtp#[        [        U5      5       H  nU R                  X#U   5      u  o#U'   M     [        [        U5      S-
  5       H  nU R                  X#U   5      u  o#U'   M     U/UQ7$ )Nr-  r   )rh  r}  r�   )r   r€  ÚheadÚtailrº  s        r   rº  Ú TritonSemantic.broadcast_tensorsã  sy   € ÞØØˆˆÜ”s˜4“yÖ!ˆAØ ×5Ñ5°dÀ¹GÓD‰MˆD�q“'ñ "ä”s˜4“y 1‘}Ö%ˆAØ ×5Ñ5°dÀ¹GÓD‰MˆD�q“'ñ &àˆ}�t‰}Ðr!   c           	     ó  ^ ^• T R                   " U6 n[        U5      S:”  d   S5       eU Vs/ s H*  oQS   R                  R                  UR                  5      PM,     nnT R
                  R                  U Vs/ s H  ofR                  PM     snU Vs/ s H  oUR                  T R
                  5      PM     snU5      mU" T5        TR                  5       (       d   e[        UU 4S j[        U5       5       5      $ s  snf s  snf s  snf )Nr   z1map_elementwise must have at least 1 input tensorc              3  ój   >#   • U  H(  u  pTR                  TR                  U5      U5      v •  M*     g 7fr.   )r+   r  )r©  rº  rõ  Úelementwise_opr   s      €€r   r«  Ú1TritonSemantic.map_elementwise.<locals>.<genexpr>û  s.   øé € ÐhÒPgÁuÀq�T—[‘[ ×!:Ñ!:¸1Ó!=¸r×BÐBÒPgùs   ƒ03)rº  r}  r{   r±   r‹   r,   Úcreate_map_elementwiser°   r²   r†  rƒ  r·  )r   r€  Úresult_typesÚpackr‡  rõ  rz  rŸ  s   `      @r   Úmap_elementwiseÚTritonSemantic.map_elementwiseí  sÝ   ù€ à×'Ò'¨Ð0ˆä�6‹{˜Q‹ÐSÐ SÓSˆÙLXÓYÊLÀb˜q™	Ÿ™×6Ñ6°r·y±yÖAÉLˆÐYØŸ™×<Ñ<Ù%Ó&šv˜!�XŒX™vÑ&Ù.:Ó;ªl¨�X‰X�d—l‘lÖ#©lÑ;Øó
ˆñ
 	˜.Ô)Ø×$Ñ$×&Ñ&Ð&Ð&äÕhÔPYÐZfÔPgÓhÓhÐhùò Zùâ&ùÚ;s   ¬1C=Á8DÂ$D
c                óþ  • [        UR                  5      S:X  d   S5       eUR                  R                  5       (       d   S5       eUb[  U R	                  X1R                  5      nUR
                  R                  R                  5       (       d  [        S5      eUR                  nU R                  U R                  R                  UR                  X#5      [        R                  " [        R                  U/5      5      $ )Nr   z histogram only supports 1D inputz%histogram only supports integer inputr¹  )r}  ri  rB   r\   r¼  r{   r‹   rç  r3   r°   r+   r,   Úcreate_histogramr5   rb  r6   )r   r¸   Únum_binsrK  s       r   Ú	histogramÚTritonSemantic.histogram  sÄ   € Ü�5—;‘;Ó 1Ó$ÐHÐ&HÓHÐ$Ø�{‰{×!Ñ!×#Ñ#ÐLÐ%LÓLÐ#ØÑØ×,Ñ,¨T·;±;Ó?ˆDØ—9‘9×#Ñ#×+Ñ+×-Ñ-Ü Ð!EÓFÐFØ—;‘;ˆDØ�{‰{˜4Ÿ<™<×8Ñ8¸¿¹ÀxÓVÜŸ=š=¬¯©°H°:Ó>ó@ð 	@r!   c                ó  • [        S[        UR                  5      5      [        U5      :w  a  [        S5      eUR                  R                  S[        R                  " X!R                  R                  5       5      5        U$ )Nr   zAShape of input to multiple_of does not match the length of valuesztt.divisibility)	rý   r}  ri  r3   r°   Úset_attrr   Ú	make_attrÚget_context©r   r|   r¿  s      r   Úmultiple_ofÚTritonSemantic.multiple_of  sY   € Üˆq”#�a—g‘g“,Ó¤3 v£;Ó.ÜÐ`ÓaÐaØ	�‰×ÑÐ+¬R¯\ª\¸&Ç(Á(×BVÑBVÓBXÓ-YÔZØˆr!   c                óð   • [        UR                  5      [        U5      :w  a  [        S5      eUR                  R	                  S[
        R                  " X!R                  R                  5       5      5        U$ )NzDShape of input to max_contiguous does not match the length of valuesztt.contiguity©r}  ri  r3   r°   r¬  r   r­  r®  r¯  s      r   Úmax_contiguousÚTritonSemantic.max_contiguous  sQ   € Üˆq�w‰w‹<œ3˜v›;Ó&ÜÐcÓdÐdØ	�‰×Ñ˜/¬2¯<ª<¸ÇÁ×@TÑ@TÓ@VÓ+WÔXØˆr!   c                óð   • [        UR                  5      [        U5      :w  a  [        S5      eUR                  R	                  S[
        R                  " X!R                  R                  5       5      5        U$ )NzCShape of input to max_constancy does not match the length of valuesztt.constancyr³  r¯  s      r   Úmax_constancyÚTritonSemantic.max_constancy  sQ   € Üˆq�w‰w‹<œ3˜v›;Ó&ÜÐbÓcÐcØ	�‰×Ñ˜.¬"¯,ª,°v¿x¹x×?SÑ?SÓ?UÓ*VÔWØˆr!   c                ór   • U R                  U R                  R                  5       [        R                  5      $ r.   )r+   r,   Úcreate_barrierr5   rq  )r   s    r   Údebug_barrierÚTritonSemantic.debug_barrier  s$   € Ø�{‰{˜4Ÿ<™<×6Ñ6Ó8¼"¿'¹'ÓBÐBr!   c                óô  • UR                  S5      (       d  U(       a  US-  nUR                  S5      (       d  U(       a  US S S-   n[        U5      S:”  a  UR                  S5      (       d  SU-   nU Vs/ s H  oDR                  PM     nnU Vs/ s H  oDR                  R                  5       PM     nnU R                  U R                  R                  XXV5      [        R                  5      $ s  snf s  snf )NÚ rµ  r—  r   )Úendswithr}  Ú
startswithr°   rB   rÓ   r+   r,   Úcreate_printr5   rq  )r   ÚprefixÚargsÚhexÚargÚnew_argsÚ	is_signeds          r   Údevice_printÚTritonSemantic.device_print"  sÉ   € ð �‰˜s×#Ñ#®Ø�c‰MˆFØ�‰˜t×$Ñ$®Ø˜C˜R�[ 4Ñ'ˆFÜˆv‹;˜‹? 6×#4Ñ#4°S×#9Ñ#9Ø˜6‘\ˆFá*.Ó/ª$ 3—J”J©$ˆÐ/Ù:>Ó?º$°3—Y‘Y×,Ñ,Ö.¹$ˆ	Ð?Ø�{‰{˜4Ÿ<™<×4Ñ4°VÀ(ÓVÔXZ×X_ÑX_Ó`Ð`ùò 0ùÚ?s   Á6C0Â#C5c                ó  • U R                   R                  R                  (       d  g Ub   U R                  XR	                  U5      5      nU R                  U R                   R                  UR                  U5      [        R                  5      $ r.   )
r,   r¡   Údebugr  r  r+   Úcreate_assertr°   r5   rq  )r   r«   r¬   rK  s       r   r¦   ÚTritonSemantic.device_assert0  s^   € Ø�|‰|×#Ñ#×)×)ØØÑØ—8‘8˜D§)¡)¨D£/Ó2ˆDØ�{‰{˜4Ÿ<™<×5Ñ5°d·k±kÀ3ÓGÌÏÉÓQÐQr!   c                óˆ   • U R                  U R                  R                  UR                  5      [        R
                  5      $ r.   )r+   r,   Úcreate_assumer°   r5   rq  )r   r«   s     r   ÚassumeÚTritonSemantic.assume7  s*   € Ø�{‰{˜4Ÿ<™<×5Ñ5°d·k±kÓBÄBÇGÁGÓLÐLr!   c                óÒ  • [        U[        5      (       a  [        R                  " U5      n[        U[        R                  5      (       añ  [        UR                  [
        5      (       a%  U R                  R                  UR                  5      $ U(       aS  SUR                  s=::  a  S:  d  O   SUR                   S35       eU R                  R                  UR                  5      $ SUR                  s=::  a  S:  d  O   SUR                   S35       eU R                  R                  UR                  5      $ [        U[        R                  5      (       aû  UR                  R                  S:X  d   S	5       eUR                  R                  5       (       d   S
5       eUR                  [        R                  :w  a^  U(       aW  U R                  R                  UR                   U R                  R#                  5       UR                  R%                  5       5      $ UR                  [        R                  :X  a  U(       d   S5       eUR                   $  S['        U5       35       e)Nrg   rh   z@Block pointers only support 64 bit `shape/strides`, got a value z which is out of the rangere   rf   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: )rn   rr   r5   ry   rR   ro   r,   rp   rî  r)  r+   rt  rB   r\   rt   r³   r°   Úget_int64_tyrÓ   r{   )r   rA  rb  s      r   Ú_convert_elem_to_ir_valueÚ(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                ó”   • [        US5      (       a!  U Vs/ s H  o0R                  X25      PM     sn$ U R                  X5      /$ s  snf )Nr<  )r=  rÔ  )r   Ú	list_likerb  rA  s       r   rd  Ú$TritonSemantic._convert_to_ir_valuesT  sI   € Ü�9˜j×)Ñ)ÙR[Ó\ÒR[È$×2Ñ2°4ÖEÑR[Ñ\Ð\Ø×.Ñ.¨yÓFÐGÐGùò ]s   –Ac           	     ó  ^• U R                  U5      nU R                  U5      nU R                  USS9nUR                  R                  5       (       a)  UR                  R                  R	                  5       (       a  [        S5      eUR                  R                  [        R                  :X  aI  U R                  U[        R                  " [        R                  UR                  R                  5      5      n[        TS5      (       d  T/mT Vs/ s H0  n[        U[        R                  5      (       a  UR                  OUPM2     snm[!        S T 5       5      (       d   S5       e[        US5      (       d  U/nU Vs/ s H0  n[        U[        R                  5      (       a  UR                  OUPM2     nn[#        U5      [%        ['        [)        U5      5      5      :X  d   S5       e[!        U4S jX#XF4 5       5      (       d   S	5       eU R*                  R-                  UR.                  X#UTU5      nU R1                  U[        R                  " [        R2                  " UR                  R                  T5      5      5      $ s  snf s  snf )
NFra  zMExpected `base` to be a pointer type (but not a block pointer type or others)r<  c              3  óz   #   • U  H1  n[        U[        5      =(       a    S Us=:*  =(       a    S:  Os  v •  M3     g7f)re   rf   N)rn   rr   )r©  rA  s     r   r«  Ú0TritonSemantic.make_block_ptr.<locals>.<genexpr>l  s.   é € Ð\ÒP[È”:˜d¤CÓ(×C¨V°t×-CÓ-C¸eÔ-CÔCÒP[ùs   ‚9;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  óR   >#   • U  H  n[        T5      [        U5      :H  v •  M     g 7fr.   )r}  )r©  r×  r@  s     €r   r«  rÛ  v  s!   øé € ÐhÒGg¸)”3�{Ó#¤s¨9£~Ö5ÒGgùs   ƒ$'zBExpected shape/strides/offsets/block_shape to have the same length)rd  r{   rƒ   rH  r|  r3   r5   rq   r�   rT  rS  rU  r=  rn   ry   rR   r„  r­  r®  rh  r}  r,   Úcreate_make_block_ptrr°   r+   rb  )	r   Úbaseri  Ústridesrg  r@  ÚorderrA  r°   s	        `   r   Úmake_block_ptrÚTritonSemantic.make_block_ptrY  s  ø€ ð ×*Ñ*¨5Ó1ˆØ×,Ñ,¨WÓ5ˆØ×,Ñ,¨WÀ%Ð,ÐHˆð �y‰y×Ñ×!Ñ! T§Y¡Y×%9Ñ%9×%BÑ%B×%DÑ%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ÒZeÐRV¤Z°´b·l±l×%CÑ%C�t—z’zÈÒMÑZeÑfˆÜÑ\ÑP[Ó\×\Ñ\ð 	VØUó	VÐ\ô �u˜j×)Ñ)Ø�GˆEÙTYÓZÒTYÈDœz¨$´·±×=Ñ=�—’À4ÒGÑTYˆÐZÜ�e‹}¤¤U¬3¨u«:Ó%6Ó 7Ó7ÐwÐ9wÓwÐ7ô ÔhÈÐX_ÑGgÓh×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   Ä7JÆ 7J
c                ó¢   • U R                  USS9nU R                  U R                  R                  UR                  U5      UR
                  5      $ ro  )rd  r+   r,   Úcreate_advancer°   r{   )r   rÞ  rg  s      r   ÚadvanceÚTritonSemantic.advance  sC   € à×,Ñ,¨WÀ%Ð,ÐHˆð �{‰{˜4Ÿ<™<×6Ñ6°t·{±{ÀGÓLÈdÏiÉiÓXÐXr!   c           	     ó<  • [        U5      nSUs=::  a  S::  d  O  [        SU S35      e[        U5      U:w  a  [        SU S[        U5       35      e[        U5      U:w  a  [        SU S[        U5       35      e[        UR                  [        R
                  5      (       d   eUR                  R                  R                  S	-  n[        R                  " US
   5      nX‡-  S:  a  [        SU SU SX‡-   S35      e[        R                  " US
   5      n	U	S:w  a  [        SU	 35      eU V
s/ s H"  o R                  U
[        R                  5      PM$     nn
U V
s/ s H6  o R                  [        R                  " U
5      [        R                  5      PM8     nn
[        R                  " U5      n[        UR                  [        R
                  5      (       d   e[        R                  " UR                  R                  U5      nUR                  nUR                  R                  R!                  5       nU R#                  U5      nUR                  R                  R%                  5       (       a)  U[&        R(                  R*                  :X  a  [        S5      eU R,                  R/                  XÂ Vs/ s H  oÿR                  PM     snU Vs/ s H  oÿR                  PM     snXMU5      n[        R0                  " UX#U5      $ s  sn
f s  sn
f s  snf s  snf )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 r¥  r—  r   zRDescriptor block shape must have at least 16 bytes in the last dimension, but got z * z = z bytesz-Tensor descriptor last dim must be 1 but got z8Padding option `nan` is not supported for integer blocks)r}  r3   rn   rB   r5   rT  rH  rÕ  rŒ  ru  r6   rt   Ú_unwrap_shaper{   rb  r°   rÓ   r  r\   r   r  r  r,   Úcreate_make_tensor_descriptorÚtensor_descriptor)r   rÞ  ri  rß  r@  r  rh  Ú	elem_sizeÚcontig_dim_sizeÚlast_strider|   r{   Úbase_handleÚis_signed_intr  rˆ  r°   s                    r   Úmake_tensor_descriptorÚ%TritonSemantic.make_tensor_descriptor†  sÌ  € ä�5‹zˆØ�T•˜Q•ÜÐ?À¸vÀ[ÐQÓRÐRÜˆw‹<˜4ÓÜ˜y¨¨Ð.?ÄÀGÃ¸~ÐNÓOÐOÜˆ{Ó˜tÓ#ÜÐ<¸T¸FÐBVÔWZÐ[bÓWcÐVdÐeÓfÐfÜ˜$Ÿ*™*¤b§o¡o×6Ñ6Ð6Ð6Ø—J‘J×)Ñ)×<Ñ<ÀÑAˆ	Ü×1Ò1°+¸b±/ÓBˆØÑ&¨Ó+ÜØdÐetÐduÐuxð  zCð  yDð  DGð  HWñ  Hcð  Gdð  djð  kóð ô ×-Ò-¨g°b©kÓ:ˆØ˜!ÓÜÐLÈ[ÈMÐZÓ[Ð[á8=Ó>º°1×!Ñ! !¤R§X¡XÖ.¹ˆÐ>ÙSZÓ[ÒSZÈa×#Ñ#¤B×$;Ò$;¸AÓ$>ÄÇÁÖIÑSZˆÐ[ô ×&Ò& {Ó3ˆä˜$Ÿ)™)¤R§_¡_×5Ñ5Ð5Ð5Ü�}Š}˜TŸY™Y×1Ñ1°;Ó?ˆØ—k‘kˆØŸ	™	×,Ñ,×:Ñ:Ó<ˆà×-Ñ-¨nÓ=ˆà�9‰9×Ñ×&Ñ&×(Ñ(¨W¼×8IÑ8I×8QÑ8QÓ-QÜÐWÓXÐXà—‘×;Ñ;¸KÐ\aÓIbÒ\aÐWXÏ(Ì(Ñ\aÑIbÙOVÓ<WÊwÈ!¿X¼XÉwÑ<WÐYdØ<CóEˆô ×#Ò# F¨E¸DÓAÐAùò' ?ùÚ[ùò JcùÚ<Ws   Ä3)L
Å"=LÊ9L
ËLr/   )r8   rr   Úreturnr   )rF   útl.dtyperG   rô  ró  rô  )rF   rô  r]   ro   rG   rô  r^   ro   r_   ro   ró  rô  )T)r}   ro   )r   rô  r   rô  r…   ro   ró  ÚNone)FFTF)r‘   úTensorTy | numbers.Numberr’   rö  ró  úTuple[TensorTy, TensorTy])r‘   r   r’   r   r§   Úcallable)r¸   rö  r¹   rö  r¢   ro   ró  r   )r¸   rö  r¹   rö  ró  r   )r¸   rö  r¹   rö  rÚ   ro   ró  r   )r|   r   rî   r   rï   útl.PropagateNan)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ó  útl.block_type)rd  rr   re  rr   rÖ   rú  ró  r   )rB   rô  ró  r   )ri  ú	List[int]rB   rô  ró  r   )rR   r   ri  rû  ró  r   )rR   r   ró  r   )r¸   r   r†  rû  r‡  ro   ró  r   )r¸   r   r8   rr   ró  r   )r‘   r   r’   r   r‡  ro   ró  r   )rš  r   r›  r   ró  r   )rš  r   ró  r÷  )r¸   r   r°  ú
Tuple[int]ró  r   )r¸   r   ri  rü  ró  r   )r‘   r   r’   r   ró  r   )rÑ  úOptional[str])r¸   r   r×  rô  ró  r   r.   )r¸   r   r×  rô  rð  rý  ró  r   )rJ  r   rK  úOptional[TensorTy]r¹   rþ  r?  r   r  rÁ  r  rÁ  r  rÁ  rL  ro   ró  r   )rf  útl.tensor_descriptor_baser  rÁ  r  rÁ  ró  r   )rf  rÿ  rR   r   ró  rõ  )rf  rÿ  rR   r   ró  r   )r  rÁ  r  rÁ  ró  r   )rJ  r   rµ  r   rK  rþ  r  rÁ  r  rÁ  ró  r   )rJ  r   rÇ  r   rµ  r   r,  rÁ  r8  rÁ  ró  r   )
rJ  r   rµ  r   rK  r   rÎ  rÁ  ró  z#Tuple[TensorTy, TensorTy, TensorTy])r|   r   ró  r   )rJ  r   rµ  r   rK  r   r,  rÁ  r8  rÁ  ró  r   )r‘   r   r’   r   r/  r   r  rý  r0  rr   r1  rô  ró  r   )rA  rÁ  )rµ  r   rA  rÁ  )r‘   r   rS  r   r\  rÁ  r’   r   rT  rþ  r]  rÁ  r/  zTensorTy | Noner^  ro   r_  ro   r`  ro   r1  rô  ró  r   )rr  r   r|   r   rî   r   ró  r   )r€  úSequence[TensorTy]r8   rr   ró  úTuple[TensorTy, ...])r€  r   r8   rr   r�  ro   ró  r  )r•  r   r–  r   r8   rr   ró  r   )r€  zSequence[tl.tensor]r¢  zSequence[tl.dtype]r£  rr   ró  zTuple[tl.tensor, ...])r¸   r   r¨  rr   rK  rþ  ró  r   )r|   r   r¿  rû  ró  r   )ró  r   )rÂ  rÁ  rÃ  úList[TensorTy]rÄ  ro   ró  r   )r«   r   r¬   rÁ  rK  rþ  ró  r   )rÞ  r   ró  r   )r  )rÞ  r   ri  r  rß  r  r@  zList[tl.constexpr]r  rÁ  ró  ztl.tensor_descriptor)r"   r#   r$   r%   r5   r+   Ú__annotations__Úlangr   r9   r=   rL   rb   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;  rB  r¥   rN  r¤   rX  r^  rk  rv   ru  ry  rx  r‚  r‰  r�  r”  rŸ  r¥  r±  r¼  r�   rÒ  r  r�   r  r	  r  r  r-  r9  rC  rN  rY  r\  ri  rl  rr  r{  r…  rˆ  rŒ  r‘  r•  rš  rž  r¬  r±  r¶  r¾  rÃ  rÈ  rÑ  rÙ  rí  rñ  rõ  rú  rÿ  r  r  r  r=  rC  rO  rX  rn  rà  rv  r‰  r‘  r—  rº  r¤  r©  r°  r´  r·  r»  rÈ  r¦   rÐ  rÔ  rd  rá  rå  rñ  r&   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ôSôfôXôVôô0

ôTôWô 2ôptôjö$i9ò^òò	ò	òòò
òò,:ðxnØ ðnØ25ðnØHKðnØZ^ðnØckônð	/Ø),ð	/Ø19ô	/ô/ônôuò@ò[ô
uôuôuôuôuô$ô0*ò*xò8.$ð`WØ"ðWØ'/ôWô(uðØ'*ðØ/Rôô8+ô$)ôL$)ôL%ôyôxôyôò
<ð[yØ#&ð[yØ3;ð[yØ@Hô[yôzô0ò$hð3tØ0ð3tØ>Að3tØHWð3tØdhð3tà#ð3tà15ð3tàBJð3tàOWô3tôreò(&ôuð,xØ"&ðxØ+?ôxô.Kò,ðiØ.Côiô*	@ôôôôCôaôRôMòYô4Hô
$fôLYð W]ð(BØ,>ð(BØPSð(BØau÷(Bð (Br!   r)   )Ú
__future__r   r$  Útypingr   r   r   r   r   r	   r
   r‰   Útriton.runtimer   Ú_C.libtritonr   r¡  r   r5   r   r   Ú	Exceptionr   r)   r-  r!   r   Ú<module>r
     sV   ðÝ "Û ç J× JÑ JÛ å !å Ý áˆCƒL€Ù�:Ó€ôF 	ô FôUB�W˜XÑ&õ UBr!   