Ë
    ÏÍ:jž�  ã            	      ó”  — d dl mZmZ d dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
mZ d dlmZ d dlmZ d dlmZmZmZmZmZmZmZmZmZmZ d dlmZ d dlmZ d	d
lm Z  d	dl!m"Z" d	dl#m$Z$m%Z%m&Z&m'Z' e(d e)d«        Z* ed«      Z+ G d„ dejX                  «      Z-d)d„Z. G d„ d«      Z/i Z0g Z1d„ Z2d*d„Z3 G d„ dee+   «      Z4d„ Z5d„ Z6d„ Z7e G d„ d«      «       Z8 G d„ de4e+   «      Z9ed+d„«       Z:edddddddd œ	 	 	 	 	 	 	 	 	 	 	 	 	 d,d!„«       Z:	 d-dddddddd œ	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d.d"„Z: G d#„ d$«      Z; G d%„ d&«      Z<d'„ Z=d(„ Z>y)/é    )ÚannotationsÚdivisionN)Údefaultdict)Ú	dataclass)Úcached_property)
ÚCallableÚGenericÚIterableÚOptionalÚTypeVarÚUnionÚoverloadÚDictÚAnyÚTuple©ÚTensorDescriptor)Ú
ModuleTypeé   )Úknobs)Údriver)Úfind_paths_ifÚget_iterable_pathÚtype_canonicalisation_dictÚcanonicalize_dtypez.runtime.jitÚTc                  óv   ‡ — e Zd ZdZdˆ fd„Zed„ «       Zd„ Zd„ Zd„ Z	d„ Z
d„ Zd	„ Zd
„ Zd„ Zd„ Zd„ Zd„ Zˆ xZS )ÚDependenciesFinderaÑ  
    This AST visitor is used to find dependencies of a JITFunction. This can
    be used to invalidate a JITFunction's hash when its source code -- or
    that of its dependencies -- changes.

    This visitor also keeps track of the global variables touched by the
    JITFunction.  When we launch the kernel, we check that these have the same
    values as they did when we ran this visitor.  If not, we raise an error (or
    otherwise we could recompile).
    c                óÌ   •— t         ‰| �  «        || _        t        j                  |j                  d«      «      | _        || _        || _        h d£| _	        i | _
        d| _        y )Núutf-8>
   ÚintÚlenÚmaxÚminÚlistÚfloatÚprintÚrangeÚgetattrÚ
isinstanceF)ÚsuperÚ__init__ÚnameÚhashlibÚsha256ÚencodeÚhasherÚglobalsÚ	nonlocalsÚsupported_python_builtinsÚused_global_valsÚvisiting_arg_default_value)Úselfr-   r2   r3   ÚsrcÚ	__class__s        €úg/home/mcse/projects/srt_converter/srt-converter-venv/lib/python3.12/site-packages/triton/runtime/jit.pyr,   zDependenciesFinder.__init__)   s\   ø€ Ü‰ÑÔØˆŒ	Ü—n‘n S§Z¡Z°Ó%8Ó9ˆŒð ˆŒØ"ˆŒò*
ˆÔ&ð. TVˆÔà*/ˆÕ'ó    c                ó6   — | j                   j                  «       S ©N)r1   Ú	hexdigest©r7   s    r:   ÚretzDependenciesFinder.retN   s   € à�{‰{×$Ñ$Ó&Ð&r;   c                ó†   — t        j                  |j                  «      ryt        |dd«      }|j	                  t
        «      S )NTÚ
__module__Ú )ÚinspectÚ	isbuiltinÚfuncr)   Ú
startswithÚTRITON_MODULE)r7   ÚnoderF   Úmodules       r:   Ú_is_triton_builtinz%DependenciesFinder._is_triton_builtinR   s6   € Ü×Ñ˜TŸY™YÔ'ØÜ˜˜|¨RÓ0ˆØ× Ñ ¤Ó/Ð/r;   c                ó>  — t        |t        «      �r| j                  j                  «       |j                  j                  «       z  D ]_  }|\  }}| j                  |   \  }}|j                  |   \  }}||k7  sŒ2t	        d|› d|› d| j
                  › d|j                  › d|› d�«      ‚ | j                  j                  |j                  «       |j                  }|t        t        |dd«      «      z  }| j                  j                  |j                  d	«      «       y y )
NúGlobal variable z has value z when compiling z, but inner kernel z has conflicting value z7 from when it was first compiled.  This is not allowed.ÚnoinlineFr    )r*   ÚJITFunctionr5   ÚkeysÚRuntimeErrorr-   Ú__name__ÚupdateÚ	cache_keyÚstrr)   r1   r0   )r7   rF   ÚkÚvar_nameÚ_Úv1Úv2Úfunc_keys           r:   Ú_update_hashzDependenciesFinder._update_hashX   s3  € Ü�dœKÕ(ð ×*Ñ*×/Ñ/Ó1°D×4IÑ4I×4NÑ4NÓ4PÑPò �Ø‘�˜!Ø×-Ñ-¨aÑ0‘��AØ×-Ñ-¨aÑ0‘��AØ˜“8Ü&Ø*¨8¨*°KÀ¸tÐCSÐTX×T]ÑT]ÐS^Ð^qÐrv×rÑrð  rAð  AXð  Y[ð  X\ð  \Sð  Tóð ðð ×!Ñ!×(Ñ(¨×)>Ñ)>Ô?à—~‘~ˆHØœœG D¨*°eÓ<Ó=Ñ=ˆHØ�K‰K×Ñ˜xŸ™¨wÓ7Õ8ð )r;   c                ó  ‡ — t        |j                  «      t        j                  u r|j                  S |j                  ‰ j
                  v ry ˆ fd„} ||j                  «      \  }}|�‹‰ j                  st        |«      t        urnt        |t        «      s^t        |dd«      sQ|j                  ‰ j                  vr9t        j                  |«      |f‰ j                  |j                  t	        |«      f<   ‰ j                  |«       |S )Nc                ó¶   •— ‰j                   j                  | d «      }|�|‰j                   fS ‰j                  j                  | d «      }|�|‰j                  fS y)N)NN)r2   Úgetr3   )r-   Úvalr7   s     €r:   Úname_lookupz2DependenciesFinder.visit_Name.<locals>.name_lookupr   sZ   ø€ Ø—,‘,×"Ñ" 4¨Ó.ˆCØˆØ˜DŸL™LÐ(Ð(Ø—.‘.×$Ñ$ T¨4Ó0ˆCØˆØ˜DŸN™NÐ*Ð*Ør;   Ú__triton_builtin__F)ÚtypeÚctxÚastÚStoreÚidÚlocal_namesr6   r   r*   rO   r)   r4   Úcopyr5   r\   )r7   rI   ra   r`   Úvar_dicts   `    r:   Ú
visit_NamezDependenciesFinder.visit_Namej   sÖ   ø€ Ü�—‘‹>œSŸY™YÑ&Ø—7‘7ˆNà�7‰7�d×&Ñ&Ñ&àô	ñ $ D§G¡GÓ,‰ˆˆXð
 ˆOð ×7Ò7ô ˜“I¤ZÑ/ô # 3¬Ô4¼WÀSÐJ^Ð`eÔ=fØ—G‘G 4×#AÑ#AÑAÜ>B¿i¹iÈ»nÈhÐ=WˆD×!Ñ! 4§7¡7¬B¨x«LÐ"9Ñ:à×Ñ˜#ÔØˆ
r;   c                ó^   — |j                   D �cg c]  }| j                  |«      ‘Œ c}S c c}w r=   )ÚeltsÚvisit)r7   rI   Úelts      r:   Úvisit_TuplezDependenciesFinder.visit_Tuple‘   s$   € ð ,0¯9©9Ö5 C�—
‘
˜3•Ò5Ð5ùÒ5s   �*c                óX  — | j                  |j                  «      }t        |t        j                  «      r6| j                  |j                  «      }t        |t        j                  «      rŒ6|�t        |dd«      t        k(  ry t        ||j                  «      }| j                  |«       |S )NrR   rC   )	rn   Úvaluer*   re   Ú	Attributer)   rH   Úattrr\   )r7   rI   Úlhsr@   s       r:   Úvisit_Attributez"DependenciesFinder.visit_Attribute–   s   € Ø�j‰j˜Ÿ™Ó$ˆÜ˜œcŸm™mÔ,Ø—*‘*˜SŸY™YÓ'ˆCô ˜œcŸm™mÕ,àˆ;œ7 3¨
°BÓ7¼=ÒHØÜ�c˜4Ÿ9™9Ó%ˆØ×Ñ˜#ÔØˆ
r;   c                ó–   — |j                   j                   D �ch c]  }|j                  ’Œ c}| _        | j                  |«       y c c}w r=   )ÚargsÚargrh   Úgeneric_visit)r7   rI   ry   s      r:   Úvisit_FunctionDefz$DependenciesFinder.visit_FunctionDef    s4   € à/3¯y©y¯~©~Ö>¨˜CŸG›GÒ>ˆÔØ×Ñ˜4Õ ùò ?s   ™Ac                óp  ‡ — ˆ fd„}t        j                  |j                  |j                  |j                  r|j                  gng |j
                  «      D ]  }‰ j                  |«       Œ  ||j                  «       |j                  �‰ j                  |j                  «        ||j                  «       y )Nc                ó�   •— 	 ‰j                   rJ ‚d‰_         | D ]  }|€Œ‰j                  |«       Œ 	 d‰_         y # d‰_         w xY w)NTF)r6   rn   )ÚdefaultsÚexprr7   s     €r:   Úvisit_defaultsz:DependenciesFinder.visit_arguments.<locals>.visit_defaults±   sS   ø€ ð8Ø×:Ò:Ð:Ð:Ø26�Ô/Ø$ò )�DØÑ'ØŸ
™
 4Õ(ñ)ð 38�Õ/ø°%�Õ/ús   ƒ<  < ¼	A)
Ú	itertoolsÚchainÚposonlyargsrx   ÚvarargÚ
kwonlyargsrn   Úkw_defaultsÚkwargr~   )r7   rI   r€   ry   s   `   r:   Úvisit_argumentsz"DependenciesFinder.visit_arguments¥   sŒ   ø€ ô	8ô —?‘? 4×#3Ñ#3°T·Y±YÐQU×Q\ÒQ\ÀÇÁÁÐbdÐfj×fuÑfuÓvò 	ˆCØ�J‰J�s�Oð	ñ 	�t×'Ñ'Ô(à�:‰:Ð!Ø�J‰J�t—z‘zÔ"á�t—}‘}Õ%r;   c                óº   — | j                  |«      }t        |t        «      r| xj                  t	        |«      z  c_        y | j                  j                  |«       y r=   )rn   r*   r%   rh   ÚsetÚadd)r7   rI   Útargets      r:   ÚvisitAssnTargetz"DependenciesFinder.visitAssnTargetÅ   sE   € ð —‘˜DÓ!ˆÜ�fœdÔ#Ø×Ò¤ F£Ñ+Öà×Ñ× Ñ  Õ(r;   c                ó¨   — t        |j                  «      dk7  rt        d«      ‚| j                  |j                  d   «       | j	                  |«       y )Né   z2Simultaneous multiple assignment is not supported.r   )r"   ÚtargetsÚ	TypeErrorr�   rz   ©r7   rI   s     r:   Úvisit_AssignzDependenciesFinder.visit_AssignÎ   sG   € Üˆt�|‰|Ó Ò!ô
 ÐPÓQÐQà×Ñ˜TŸ\™\¨!™_Ô-ð 	×Ñ˜4Õ r;   c                ó\   — | j                  |j                  «       | j                  |«       y r=   ©r�   rŒ   rz   r’   s     r:   Úvisit_AnnAssignz"DependenciesFinder.visit_AnnAssignÛ   ó$   € Ø×Ñ˜TŸ[™[Ô)ð 	×Ñ˜4Õ r;   c                ó\   — | j                  |j                  «       | j                  |«       y r=   r•   r’   s     r:   Ú	visit_ForzDependenciesFinder.visit_Forá   r—   r;   )ÚreturnÚNone)rR   rB   Ú__qualname__Ú__doc__r,   Úpropertyr@   rK   r\   rk   rp   rv   r{   rˆ   r�   r“   r–   r™   Ú__classcell__©r9   s   @r:   r   r      s[   ø„ ñ	õ#0ðJ ñ'ó ð'ò0ò9ò$%òN6ò
ò!ò
&ò@)ò!ò!ö!r;   r   c                ó   — dd l mc m} t        | t        «      rÇ| j                  «       } | j                  d«      r7| j                  d«      } t        | «      } | j                  d«      sJ ‚d| dd  z   S | j                  d«      rdt        | d d «      z   S | j                  d«      rdt        | dd  «      z   S | j                  d«      r“t        | j                  d«      «      S t        | |j                  «      rdt        | j                  «      › �S t        | |j                  «      r| j                  } n(t        | t        «      r| j                  } nt	        | «      } t!        j"                  | j%                  d	d
«      | «      S )Nr   zconst ÚconstÚ*ú*kr�   éÿÿÿÿztl.Ú_trC   )Útriton.language.coreÚlanguageÚcorer*   rU   ÚstriprG   ÚremoveprefixÚ_normalize_tyÚendswithÚpointer_typeÚ
element_tyÚdtyper-   rc   rR   r   r_   Úreplace)Útyr©   s     r:   r¬   r¬   í   s?  € ß'Ð'Ü�"”cÔØ�X‰X‹ZˆØ�=‰=˜Ô"Ø—‘ Ó)ˆBÜ˜rÓ"ˆBØ—=‘= Ô%Ð%Ð%Ø˜"˜Q˜R˜&‘=Ð Ø�;‰;�sÔØœ r¨#¨2 wÓ/Ñ/Ð/Ø�=‰=˜ÔØœ r¨!¨" vÓ.Ñ.Ð.Ø�=‰=˜ÔÜ  §¡°Ó!7Ó8Ð8Ü	�B˜×)Ñ)Ô	*Ø”= §¡Ó/Ð0Ð1Ð1Ü	�B˜Ÿ
™
Ô	#Ø�W‰W‰Ü	�BœÔ	Ø�[‰[‰ä�‹WˆÜ%×)Ñ)¨"¯*©*°T¸2Ó*>ÀÓCÐCr;   c                  ó�   — e Zd ZdZ	 	 dd„Zed„ «       Zedd„«       Zedd„«       Zed„ «       Z	ed„ «       Z
ed„ «       Zed	„ «       Zy
)ÚKernelParamzBRepresents a parameter (name plus metadata) to a @jit'ed function.c                ó<   — || _         || _        || _        || _        y r=   )ÚnumÚ_paramÚdo_not_specializeÚdo_not_specialize_on_alignment)r7   r¶   Úparamr¸   r¹   s        r:   r,   zKernelParam.__init__
  s    € àˆŒØˆŒØ!2ˆÔØ.LˆÕ+r;   c                ó.   — | j                   j                  S r=   )r·   r-   r?   s    r:   r-   zKernelParam.name  s   € à�{‰{×ÑÐr;   c                óÐ   — | j                   j                  r1| j                   j                  t        j                  j                  k(  ryt        | j                   j                  «      S )NrC   )r·   Ú
annotationrD   Ú	ParameterÚemptyr¬   r?   s    r:   r½   zKernelParam.annotation  sD   € à�{‰{×%Ò%¨¯©×)?Ñ)?Ä7×CTÑCT×CZÑCZÒ)ZØÜ˜TŸ[™[×3Ñ3Ó4Ð4r;   c                óÌ   — | j                   }|j                  d«      r|dd  }n|j                  d«      r|dd  }|t        t        j                  «       «      v r| j                   S y)Nr¤   r   r£   r�   rC   )r½   rG   rŠ   r   Úvalues)r7   Úas     r:   Úannotation_typezKernelParam.annotation_type  s]   € à�O‰OˆØ�<‰<˜ÔØ�!�"�‰AØ�\‰\˜#ÔØ�!�"�ˆAØ”Ô.×5Ñ5Ó7Ó8Ñ8Ø—?‘?Ð"Ør;   c                ó   — d| j                   v S ©NÚ	constexpr)r½   r?   s    r:   Úis_constexprzKernelParam.is_constexpr&  s   € à˜dŸo™oÐ-Ð-r;   c                ór   — | j                   ryd| j                  v xs | j                  j                  d«      S )NFr¢   r¤   )rÇ   r½   rG   r?   s    r:   Úis_constzKernelParam.is_const*  s1   € à×ÒØØ˜$Ÿ/™/Ð)ÒM¨T¯_©_×-GÑ-GÈÓ-MÐMr;   c                ó.   — | j                   j                  S r=   )r·   Údefaultr?   s    r:   rË   zKernelParam.default0  s   € à�{‰{×"Ñ"Ð"r;   c                ód   — | j                   j                  t        j                  j                  k7  S r=   )r·   rË   rD   r¾   r¿   r?   s    r:   Úhas_defaultzKernelParam.has_default4  s#   € à�{‰{×"Ñ"¤g×&7Ñ&7×&=Ñ&=Ñ=Ð=r;   N)r¶   r!   rº   zinspect.Parameterr¸   Úboolr¹   rÎ   ©rš   rU   )rR   rB   rœ   r�   r,   r   r-   r½   rÃ   rÇ   rÉ   rž   rË   rÍ   © r;   r:   r´   r´     s¢   „ ÙLðMØ15óMð ñ ó ð ð ò5ó ð5ð
 òó ðð ñ.ó ð.ð ñNó ðNð
 ñ#ó ð#ð ñ>ó ñ>r;   r´   c                ó8   ‡ ‡‡‡— ddl mŠ ddlmŠ dˆˆˆ ˆfd„	Š‰S )Nr   ©rÆ   r   r   c                óÎ  •‡ — ‰ €yt        ‰ t        «      ryt        ‰ t        «      r7|r ‰‰ d|¬«      nd }‰ dk(  r|ryd‰ k  r	‰ dk  rd	|fS d
‰ k  r	‰ dk  rd|fS d|fS t        ‰ t        «      ryt	        ‰ d«      rZ‰ j
                  |f}t        j                  |d «      }|€!|d   rdndt        |d   «      z   }|t        |<   |r ‰‰ d|¬«      nd }||fS t        ‰ t        «      rd‰ j                  fS t        ‰ ‰«      rd‰ fS t	        ‰ d«      ryt        ‰ t        «      rP‰ D �cg c]
  } ‰|«      ‘Œ }}ˆ fd„}	 |	|D �cg c]  }|d   ‘Œ	 c}«      }
 |	|D �cg c]  }|d   ‘Œ	 c}«      }|
|fS t        ‰ t        «      rTt	        ‰ j                  d«      sJ ‚t        ‰ j                  j
                  «      }d|› t        ‰ j                  «      › d�d fS t        ‰ ‰«      rat	        ‰ j                  d«      sJ ‚t        ‰ j                  j
                  «      }d|› t        ‰ j                  «      › d‰ j                   ›d�d fS t#        dt%        ‰ «      z  «      ‚c c}w c c}w c c}w )N)rÆ   N)Úu1Nr!   )Úalignr�   )rÆ   r�   i   €iÿÿÿÚi32l            l   ÿÿÿÿ Úu64Úi64)Úfp32NÚdata_ptrr¤   r£   r   ÚtensorrÆ   Útma_desc_cpu_ptr)Ú	nvTmaDescNc                óN   •— t        ‰d«      r t        ‰«      | Ž S t        | «      S )NÚ_fields)Úhasattrrc   Útuple)Úvalsry   s    €r:   ú<lambda>zAcreate_specialize_impl.<locals>.specialize_impl.<locals>.<lambda>d  s&   ø€ ¼'À#ÀyÔ:Q i¤d¨3£i°Ð&6€ ÔW\Ð]aÓWb€ r;   ztensordesc<ú>ú,zUnsupported type: %s)r*   rÎ   r!   r&   rà   r°   Ú	dtype2strr_   r   rO   rT   rá   r   Úbaser%   Úblock_shapeÚlayoutr‘   rc   )ry   rÉ   Úspecialize_valuerÕ   ÚkeyÚdskÚresÚxÚspecÚ
make_tupleÚtysrP   ÚinnerÚGluonTensorDescriptorrÆ   Úspecialize_extraÚspecialize_impls   `            €€€€r:   rõ   z/create_specialize_impl.<locals>.specialize_implB  sj  ù€ Øˆ;Ø&Ü˜œTÔ"ØÜ˜œSÔ!Ù?OÑ" 3¨°UÕ;ÐUYˆCØ�aŠxÑ,Ø'Ø˜S’ S¨IÒ%5Ø˜s�|Ð#Ø˜#’ #¨Ò"2Ø˜s�|Ð#à˜s�|Ð#Ü˜œUÔ#Ø!Ü�S˜*Ô%à—9‘9˜hÐ'ˆCÜ—-‘-  TÓ*ˆCØˆ{Ø" 1šv‘t¨3Ô2DÀSÈÁVÓ2LÑL�Ø!$”	˜#‘ÙBRÑ" 3¨¸Õ>ÐX\ˆCØ˜�:ÐÜ˜œ[Ô)Ø §¡Ð/Ð/Ü˜˜YÔ'Ø Ð%Ð%Ü�SÐ,Ô-Ø&Ü˜œUÔ#Ø03Ö4¨1‘O AÕ&Ð4ˆDÐ4ÛbˆJÙ¨DÖ1 q˜a ›dÒ1Ó2ˆCÙ¨TÖ2¨˜q ›tÒ2Ó3ˆDØ˜�;ÐÜ˜Ô-Ô.Ü˜3Ÿ8™8 ZÔ0Ð0Ð0Ü& s§x¡x§~¡~Ó6ˆEØ! % ¬¨c¯o©oÓ)>Ð(?¸qÐAÀ4ÐHÐHÜ˜Ð2Ô3Ü˜3Ÿ8™8 ZÔ0Ð0Ð0Ü& s§x¡x§~¡~Ó6ˆEØ! % ¬¨c¯o©oÓ)>Ð(?¸qÀÇÁÀÈaÐPÐRVÐWÐWäÐ2´T¸#³YÑ>Ó?Ð?ùò 5ùâ1ùÚ2s   Ä$IÅ IÅI")FTT)r¨   rÆ   Ú'triton.experimental.gluon.nvidia.hopperr   )rô   ró   rÆ   rõ   s   `@@@r:   Úcreate_specialize_implr÷   =  s   û€ å$Ýa÷/@ð /@ðb Ðr;   c                ó�   — t        t        «      dk(  rt        j                  t        d„ «      «       t        d   } || |¬«      d   S )Nr   c                 ó   — y r=   rÐ   )rX   Úkwargss     r:   rã   zmangle_type.<locals>.<lambda>x  s   � r;   )rê   )r"   Úspecialize_impl_cacheÚappendr÷   )ry   Ú
specializerõ   s      r:   Úmangle_typerþ   v  sA   € Ü
Ô Ó! QÒ&Ü×$Ñ$Ô%;Ñ<TÓ%UÔVÜ+¨AÑ.€OÙ˜3°Ô<¸QÑ?Ð?r;   c                  ó    — e Zd ZU ded<   dd„Zy)ÚKernelInterfacer   Úrunc                ó   ‡ ‡— ˆˆ fd„S )z­
        A JIT function is launched with: fn[grid](*args, **kwargs).
        Hence JITFunction.__getitem__ returns a callable proxy that
        memorizes the grid.
        c                 ó.   •—  ‰j                   | ‰ddœ|¤ŽS )NF©ÚgridÚwarmup)r  )rx   rú   r  r7   s     €€r:   rã   z-KernelInterface.__getitem__.<locals>.<lambda>†  s   ø€  x t§x¡xÈ$°TÀ%Ñ'YÐRXÑ'Y€ r;   rÐ   )r7   r  s   ``r:   Ú__getitem__zKernelInterface.__getitem__€  s   ù€ ô ZÐYr;   N)rš   r   )rR   rB   rœ   Ú__annotations__r  rÐ   r;   r:   r   r   }  s   … Ø	
ƒFôZr;   r   c           
     óø  — |j                  «       D ��ci c],  \  }}||j                  j                  dk(  rt        |«      n|“Œ. }}}dd l}| ||j                  «       D �cg c]  }t        |«      ‘Œ c}t        |j                  «       «      |j                  «       D �cg c]  }t        |«      ‘Œ c}t        |j                  «       «      |j                  dœ}	|j                  |	«      }
|
S c c}}w c c}w c c}w )Nr°   r   )r-   Ú	signatureÚconstant_keysÚconstant_valsÚ
attrs_keysÚ
attrs_valsÚoptionsrë   )
Úitemsr9   rR   rU   ÚjsonrP   r%   rÁ   Ú__dict__Údumps)r-   r
  Ú	constantsÚattrsr  rë   rr   r  rî   ÚobjÚserialized_objs              r:   Úserialize_specialization_datar  Š  sÕ   € Øen×etÑetÓev×wÑWaÐWZÐ\a� E§O¡O×$<Ñ$<ÀÒ$G”c˜%”jÈUÑRÐw€IÑwÛà 9ÐQZ×Q_ÑQ_ÓQaÖ?bÈAÄÀQÅÒ?bÜˆY×ÑÓÓ À%Ç*Á*Ã,Ö0O¸Q´°aµÒ0OÔ_cÐdi×dpÑdpÓdrÓ_sØ×#Ñ#¨Cñ€Cð
 —Z‘Z “_€NØÐùó xùò @cùÚ0Os   ”1C,Á C2ÂC7c                ó  — t        | j                  «      t        |«      k(  sJ ‚g }t        | j                  j                  «       |«      D �]  \  }}|j                  r|j                  d|› d�«       Œ)|j                  rdnd}|j                  rdnd}|j                  rdnd}d|› d|› d|› d|› d�	}	|j                  r�t        |j                  t        «      r"|j                  dk(  s|j                  dd	 d
v rd}|r#|j                  d|j                  › d|	› d�«       ŒÕ|j                  d|j                  › d�«       Œõ|j                  |	› «       �Œ	 d„ }
ddj                  t        t        |
| j                  j                  «       «      «      dgz   «      › ddj                  | j                  j                  «       D �cg c]
  }d|› d|› �‘Œ c}«      › ddj                  |«      › d�}| j                  j                  «       D ��ci c];  \  }}|j                   t"        j$                  j&                  urd|› �|j                   “Œ= }}}t(        |d<   t+        |j,                  «      |d<   t/        ||«       |d   S c c}w c c}}w )a2  
    Equivalent to sig.bind followed by apply_defaults. This generates a
    native Python function (using exec) which can be memoized on a per-kernel
    basis to avoid having to run these expensive functions -- which constitute
    much of the kernel launch overhead -- every time we run the kernel.
    z("constexpr", ú)ÚTrueÚFalsezspecialize_impl(ú, rÔ   Nr   )ÚfpÚbfFz("z",) + z[1:]z", None)c                óx   — | d   j                   t        j                  j                  u r| d   S | d   › d| d   › �S )Nr�   r   z	=default_)rË   rD   r¾   r¿   )rî   s    r:   rã   z0create_function_from_signature.<locals>.<lambda>·  sA   € ˜A˜a™DŸL™L¬G×,=Ñ,=×,CÑ,CÑC�A�a‘D€ ÈAÈaÉDÈ6ÐQZÐ[\Ð]^Ñ[_ÐZ`ÐIa€ r;   z
def dynamic_func(z	**optionsz):
    params = {ú'z': z}
    specialization = [rå   z-]
    return params, specialization, options
Údefault_rO   rõ   Údynamic_func)r"   Ú
parametersÚziprP   rÇ   rü   rÉ   r¸   r¹   rÃ   r*   rU   Újoinr%   Úmapr  rË   rD   r¾   r¿   rO   r÷   Úget_arg_specializationÚexec)ÚsigÚkparamsÚbackendÚspecializationr-   ÚkprÉ   rý   rÕ   r@   ry   Ú	func_bodyrº   Úfunc_namespaces                 r:   Úcreate_function_from_signaturer1  –  s€  € ô ˆs�~‰~Ó¤# g£,Ò.Ð.Ð.à€Nä˜Ÿ™×+Ñ+Ó-¨wÓ7ó 0‰ˆˆbØ�?Š?Ø×!Ñ! N°4°&¸Ð":Õ;à!#§¢‘v°'ˆHØ$&×$8Ò$8™¸fˆJØ!×@Ò@‘GÀfˆEØ$ T F¨"¨X¨J°b¸¸ÀBÀuÀgÈQÐOˆCØ×!Ò!Ü˜b×0Ñ0´#Ô6Ø×)Ñ)¨TÒ1°R×5GÑ5GÈÈÐ5KÈ|Ñ5[à%*˜
ÙØ"×)Ñ)¨B¨r×/AÑ/AÐ.BÀ&ÈÈÈTÐ*RÕSð #×)Ñ)¨B¨r×/AÑ/AÐ.BÀ(Ð*KÕLà×%Ñ%¨¨Ö/ð'0ñ, b€CðØ—)‘)œD¤ S¨#¯.©.×*>Ñ*>Ó*@Ó!AÓBÀkÀ]ÑRÓSÐTð UØ—	‘	¸3¿>¹>×;NÑ;NÓ;PÖQ°4˜Q˜t˜f C¨ vÒ.ÒQÓRÐSð TØ—x‘x Ó/Ð0ð 1ð€Ið Ÿ>™>×/Ñ/Ó1÷áˆD�%Ø�=‰=¤× 1Ñ 1× 7Ñ 7Ñ7ð �4�&Ð˜5Ÿ=™=Ñ(ð€Nñ ô %0€N�=Ñ!Ü(>¸w×?]Ñ?]Ó(^€NÐ$Ñ%ô 	ˆ�NÔ#ð ˜.Ñ)Ð)ùò% Rùó
s   Ç I=ÈA Jc                ó8   — | j                   › d| j                  › �S )Nú.)rB   rœ   ©Úfns    r:   Úget_full_namer6  Ð  s   € Ø�m‰mˆ_˜A˜bŸo™oÐ.Ð/Ð/r;   c                  ó,   — e Zd ZU ded<   ded<   ded<   y)ÚJitFunctionInfor   rJ   rU   r-   rO   Újit_functionN)rR   rB   rœ   r  rÐ   r;   r:   r8  r8  Ô  s   … àÓØ
ƒIØÔr;   r8  c                  ó¢   ‡ — e Zd Zd„ Z	 	 dd„Zd„ Zd„ Zd„ Zd„ Z	 	 dd„Z	d„ Z
ed	„ «       Zed
„ «       Zd„ Zd„ Zd„ Zd„ Zˆ fd„Zˆ fd„Zd„ Zˆ xZS )rO   c                 ó   — y)NFrÐ   r?   s    r:   Úis_gluonzJITFunction.is_gluonÝ  s   € Ør;   c	                ó²  — |sy | j                   j                  }	| j                   j                  }
dj                  t	        | j
                  |d   «      D ��cg c]  \  }}|j                  › d|› �‘Œ c}}«      }|	› d|j                  › d|j                  › d|j                  › d|j                  › d|j                  › d	|› d
�}t        | j                   «      }t        ||||d   ||«      }||||j                  |j                  |j                  |j                  |j                  |j                  |||dœ} |||t        |
|	| «      d|i|¥|d¬«      S c c}}w )Nr  r�   z: z[num_warps=z, num_ctas=z, num_stages=z, enable_fp_fusion=z, launch_cooperative_grid=ú](r  r   )r
  Údevicer  Ú	num_warpsÚnum_ctasÚ
num_stagesÚenable_fp_fusionÚlaunch_cooperative_gridÚextern_libsÚconfigsÚspecialization_dataÚ	is_warmuprë   F)rë   Úreprr5  ÚcompileÚis_manual_warmupÚalready_compiled)r5  rœ   rB   r&  r%  Úparamsr-   r@  rA  rB  rC  rD  r6  r  rE  r8  )r7   Úhookrë   r
  r?  r  r  rF  rH  r-   rJ   rº   r²   Ú	arg_reprsrI  Ú	full_namerG  rú   s                     r:   Ú
_call_hookzJITFunction._call_hookà  s¶  € ñ Øà�w‰w×#Ñ#ˆØ—‘×#Ñ#ˆØ—I‘IÄcÈ$Ï+É+ÐWZÐ[\ÑW]ÓF^×_¹¸À %§*¡* ¨R°¨tÒ4Ó_Ó`ˆ	Ø��{ 7×#4Ñ#4Ð"5°[À×AQÑAQÐ@RÐR_Ð`g×`rÑ`rÐ_sð  tGð  HO÷  H`ñ  H`ð  Gað  a{ð  |C÷  |[ñ  |[ð  {\ð  \^ð  _hð  ^ið  ijð  kˆÜ! $§'¡'Ó*ˆ	ä;¸IÀyÐR[Ð]dÐefÑ]gÐipÐruÓvÐð #ØØ"Ø ×*Ñ*Ø×(Ñ(Ø!×,Ñ,Ø '× 8Ñ 8Ø'.×'FÑ'FØ"×.Ñ.ØØ#6Ø"ñ
ˆñ ØØÜ˜v t¨TÓ2Ø˜CÐ* 6Ð*Ø&Ø"ô
ð 	
ùó+ `s   ÁE
c                óT   — t        |«      sJ ‚| j                  j                  |«       y)z�
        Add a hook that will be executed prior to the execution of run
        function with args and kwargs passed into the kernel
        N)ÚcallableÚpre_run_hooksrü   )r7   rN  s     r:   Úadd_pre_run_hookzJITFunction.add_pre_run_hook  s$   € ô
 ˜Œ~Ðˆ~Ø×Ñ×!Ñ! $Õ'r;   c                óÞ   — ddl m}m}m}m} t
        j                  j                  «       } ||«      }|| _        || _        || _        t        | j                  | j                  |«      }i |||fS )z1
        Precompute as much as possible.
        r   )ÚCompiledKernelrJ  Ú	ASTSourceÚmake_backend)ÚcompilerrW  rJ  rX  rY  r   ÚactiveÚget_current_targetr1  r
  rM  )r7   rW  rJ  rX  rY  rŒ   r,  Úbinders           r:   Úcreate_binderzJITFunction.create_binder  sb   € ÷ 	PÓOÜ—‘×1Ñ1Ó3ˆÙ˜vÓ&ˆØ,ˆÔØˆŒØ"ˆŒÜ/°·±ÀÇÁÈWÓUˆØ�6˜7 FÐ*Ð*r;   c               óô  — |j                  d| j                  «      xs t        j                  j                  |d<   t        j
                  j                  «       }t        j
                  j                  |«      }| j                  D ]
  } ||i |¤Ž Œ | j                  |   \  }}	}
} ||i |¤Ž\  }}}t        |«      t        |«      z   }|j                  |d «      }|�€Î|
j                  |«      }| j                  D �cg c]  }|j                  ‘Œ }}|D �cg c]  }|d   ‘Œ	 }}t        ||«      D ��ci c]  \  }}||“Œ
 }}}d|vsJ d«       ‚d|vsJ d«       ‚d|vsJ d«       ‚|D ]#  }||j                  vsŒ||vsŒt!        d	|z  «      ‚ t#        |d
„ «      }|D �ci c]&  }|t%        t'        |j)                  «       «      |«      “Œ( }}|D �cg c]  }|d   ‘Œ	 }}t#        |d„ «      }|D �ci c]  }||
j+                  t%        ||«      «      “Œ  }}| j-                  t        j                  j.                  ||||||g|«      ry | j1                  | |||«      }| j3                  ||	|j                  ¬«      }|||<   | j-                  t        j                  j4                  ||||||g|«       t7        «       }| j8                  j;                  «       D ]6  \  \  }}\  }} | j                  ||«      x}!|k7  sŒ$t=        d|› d|› d|!› �«      ‚ |sÎ|€J ‚t?        |«      r ||«      }tA        |«      }"|d   }#|"dkD  r|d   nd}$|"dkD  r|d   nd}% |jB                  ||g|j)                  «       ¢­Ž }& |jD                  |#|$|%||jF                  |jH                  |&t        j                  jJ                  t        j                  jL                  g	|j)                  «       ¢­Ž  |S c c}w c c}w c c}}w c c}w c c}w c c}w )NÚdebugr   Údevice_typez=device_type option is deprecated; current target will be usedr?  z8device option is deprecated; current device will be usedÚstreamz8stream option is deprecated; current stream will be usedz2Keyword argument %s was specified but unrecognisedc                ó   — |dk(  S rÅ   rÐ   )rX   r`   s     r:   rã   z!JITFunction.run.<locals>.<lambda>G  s   € ¸sÀkÑ?Q€ r;   r�   c                ó"   — t        |t        «      S r=   )r*   rU   )rX   rî   s     r:   rã   z!JITFunction.run.<locals>.<lambda>K  s   € ¼ÀAÄsÓ9K€ r;   )rŒ   r  rM   z1 has changed since we compiled this kernel, from z to r   )'r_   r`  r   Úruntimer   r[  Úget_current_deviceÚget_current_streamrT  Údevice_cachesrU   Úparse_optionsrM  r-   r%  r  ÚKeyErrorr   r   r%   rÁ   Ú
parse_attrrQ  Újit_cache_hookrX  rJ  Újit_post_compile_hookÚobjectr5   r  rQ   rS  r"   Úlaunch_metadatar  ÚfunctionÚpacked_metadataÚlaunch_enter_hookÚlaunch_exit_hook)'r7   r  r  rx   rú   r?  rb  rN  Úkernel_cacherŒ   r,  r]  Ú
bound_argsr-  r  rë   Úkernelrî   ÚsigkeysÚsigvalsrV   Úvr
  Ú
constexprsÚpathÚattrvalsr  r8   Únot_presentr-   rX   r`   Úglobals_dictÚnewValÚ	grid_sizeÚgrid_0Úgrid_1Úgrid_2ro  s'                                          r:   r  zJITFunction.run#  s5  € Ø Ÿ*™* W¨d¯j©jÓ9ÒP¼U¿]¹]×=PÑ=Pˆˆw‰ô —‘×1Ñ1Ó3ˆÜ—‘×1Ñ1°&Ó9ˆð ×&Ñ&ò 	"ˆDÙ�$Ð!˜&Ó!ð	"ð 15×0BÑ0BÀ6Ñ0JÑ-ˆ�f˜g vñ /5°dÐ.E¸fÑ.EÑ+ˆ
�N Gô �.Ó!¤C¨£LÑ0ˆØ×!Ñ! # tÓ,ˆð ‰>à×+Ñ+¨FÓ3ˆGà'+§{¡{Ö3 !�q—v“vÐ3ˆGÐ3Ø%3Ö4 �q˜“tÐ4ˆGÐ4Ü,/°¸Ó,A×B¡& 1 a˜˜A™ÐBˆIÑBà ¨Ñ.ÐoÐ0oÓoÐ.Ø 6Ñ)ÐeÐ+eÓeÐ)Ø 6Ñ)ÐeÐ+eÓeÐ)Øò ]�Ø˜G×,Ñ,Ò,°¸'Ò1AÜ"Ð#WÐZ[Ñ#[Ó\Ð\ð]ô ' wÑ0QÓRˆJØ_iÖjÐW[˜$Ô 1´$°z×7HÑ7HÓ7JÓ2KÈTÓ RÑRÐjˆJÐjà&4Ö5 ˜˜!›Ð5ˆHÐ5Ü! (Ñ,KÓLˆEØTYÖZÈq�Q˜×*Ñ*Ô+<¸XÀqÓ+IÓJÑJÐZˆEÐZØ�‰œuŸ}™}×;Ñ;¸SÀ)ÈVÐU_ÐahÐkpÐjqØ%ô'àà—.‘.  y°*¸eÓDˆCØ—\‘\ #¨f¸g×>NÑ>N�\ÓOˆFØ &ˆL˜ÑØ�O‰OœEŸM™M×?Ñ?ÀÀiÐQWÐYcÐelÐotÐnuØ"ô$ô “hˆØ.2×.CÑ.C×.IÑ.IÓ.Kò 	qÑ*‰IˆT�1Ñ*˜˜\Ø&×*Ñ*¨4°Ó=Ð=�À#ÓEÜ"Ø& t fÐ,]Ð^aÐ]bÐbfÐgmÐfnÐoóqð qð	qñ
 àÐ#Ð#Ð#Ü˜Œ~Ù˜JÓ'�Ü˜D›	ˆIØ˜!‘WˆFØ )¨A¢�T˜!’W°1ˆFØ )¨A¢�T˜!’W°1ˆFà4˜f×4Ñ4°T¸6ÐXÀJ×DUÑDUÓDWÒXˆOØˆF�J‰J�v˜v v¨v°v·±È×H^ÑH^Ð`oÜ—}‘}×6Ñ6¼¿¹×8VÑ8VðnØYc×YjÑYjÓYlónàˆùò_ 4ùÚ4ùÛBùò kùâ5ùâZs$   Ã=OÄO Ä3O%Æ+O+ÇO0Ç-#O5c                óT   — | j                   €| j                  S | j                  |«      S r=   )Ú_reprÚ_fn_name)r7   rX   s     r:   rI  zJITFunction.reprm  s"   € Ø $§
¡
Ð 2ˆt�}‰}ÐE¸¿
¹
À1»ÐEr;   c	           	     ó  — |r|ng }|r|ng }|| _         |j                  | _        || _        t	        j
                  |«      | _        || _        || _        t	        j                  |«      d   | _	        || _
        t        |«      | _        || _        g | _        t        | j
                  j                   j#                  «       «      D ]T  \  }	}
|	|v xs |
j$                  |v }|	|v xs |
j$                  |v }| j                  j'                  t)        |	|
||«      «       ŒV t+        j,                  t	        j.                  |«      «      }|t1        j2                  d|t0        j4                  «      j7                  «       d  }| j9                  |«       t;        | j<                  «      | _        d | _         i | _!        d | _"        || _#        || _$        | j                  D �cg c]  }|j$                  ‘Œ c}| _%        | j                  D �cg c]  }|jL                  sŒ|jN                  ‘Œ c}| _(        g | _)        |jT                  | _*        |jV                  | _+        |jX                  | _,        |jZ                  | _-        |j                  | _        y c c}w c c}w )Nr�   z^def\s+\w+\s*\().r5  rB   rJ   ÚversionrD   r
  r¸   r¹   ÚgetsourcelinesÚstarting_line_numberr…  r6  r†  ro  rM  Ú	enumerater$  rÁ   r-   rü   r´   ÚtextwrapÚdedentÚ	getsourceÚreÚsearchÚ	MULTILINEÚstartÚ_unsafe_update_srcr   r^  rh  Úhashr5   rv  r`  rN   Ú	arg_namesrÇ   r¶   rz  rT  r�   rR   rœ   Ú__globals__)r7   r5  rˆ  r¸   r¹   r`  rN   rI  ro  Úirº   ÚdnsÚdns_oar8   Úps                  r:   r,   zJITFunction.__init__p  s#  € á1BÑ-ÈÐÙKiÑ)GÐoqÐ&àˆŒØ—m‘mˆŒØˆŒÜ ×*Ñ*¨2Ó.ˆŒØ!2ˆÔØ.LˆÔ+Ü$+×$:Ñ$:¸2Ó$>¸qÑ$AˆÔ!ØˆŒ
Ü% bÓ)ˆŒØ.ˆÔàˆŒÜ! $§.¡.×";Ñ";×"BÑ"BÓ"DÓEò 	C‰HˆAˆuØÐ(Ð(ÒK¨E¯J©JÐ:KÐ,KˆCØÐ8Ð8Òh¸E¿J¹JÐJhÐ<hˆFØ�K‰K×Ñœ{¨1¨e°S¸&ÓAÕBð	Cô �o‰oœg×/Ñ/°Ó3Ó4ˆØ”"—)‘)Ð.°´R·\±\ÓB×HÑHÓJÐKÐLˆØ×Ñ Ô$ä(¨×);Ñ);Ó<ˆÔØˆŒ	ð TVˆÔð ˆŒØˆŒ
Ø ˆŒð +/¯+©+Ö6 Q˜!Ÿ&›&Ò6ˆŒØ*.¯+©+ÖH Q¸¿»˜1Ÿ5›5ÒHˆŒð  ˆÔð —z‘zˆŒØŸ™ˆŒØŸO™OˆÔØŸ>™>ˆÔØŸ-™-ˆ�ùò 7ùÚHs   ÇJÇ=JÈJc                ón   — | j                   t        j                  | j                  «      j                  z  S r=   )r–  rD   Úgetclosurevarsr5  r3   r?   s    r:   Úget_capture_scopezJITFunction.get_capture_scope®  s(   € Ø×Ñ¤'×"8Ñ"8¸¿¹Ó"A×"KÑ"KÑKÐKr;   c                óÌ  — | j                   €Ít        j                  | j                  «      j                  }t        | j                  | j                  || j                  ¬«      }|j                  | j                  «       «       |j                  t        | j                  «      z   | _         t        t        |j                   j#                  «       «      «      | _        | j                   S )N)r-   r2   r3   r8   )r”  rD   rœ  r5  r3   r   r†  r–  r8   rn   Úparser@   rU   rŠ  ÚdictÚsortedr5   r  )r7   r3   Údependencies_finders      r:   rT   zJITFunction.cache_key±  sª   € ð �9‰9ÐÜ×.Ñ.¨t¯w©wÓ7×AÑAˆIÜ"4¸$¿-¹-ÐQU×QaÑQaÐmvØ9=¿¹ô#CÐà×%Ñ% d§j¡j£lÔ3Ø+×/Ñ/´#°d×6OÑ6OÓ2PÑPˆDŒIÜ$(¬Ð0C×0TÑ0T×0ZÑ0ZÓ0\Ó)]Ó$^ˆDÔ!Ø�y‰yÐr;   c                ó   — ddl m} |S )Nr   rÒ   )r§   rÆ   )r7   rÆ   s     r:   rc   zJITFunction.type½  s   € å2ØÐr;   c               ó\   —  | j                   t        t        j                  |«      |ddœ|¤ŽS )NTr  )r  r'  Ú
MockTensorÚ
wrap_dtype)r7   r  rx   rú   s       r:   r  zJITFunction.warmupÂ  s*   € Øˆt�x‰x´´Z×5JÑ5JÈDÓ1Q˜T¨$Ñ\ÐU[Ñ\Ð\r;   c           	     ó8  — ddl m}m} dd l}dd lm} t        j                  j                  «       }|j                  |«      }|d   | j                  k7  rt        d|d   › d| j                  › �«      ‚t        t        |d   «      }|d   }	t        ||	«      D �
�ci c]4  \  }
}|
|j                  j!                  |«      r|j                  |«      n|“Œ6 }}
}t        t        |d	   «      }|d
   }t#        t        ||«      «      }t#        |d   j%                  «       «      } || |||«      }|d   j%                  «       D �
�ci c]#  \  }
}|
t'        |t(        «      rt        |«      n|“Œ% }}
}|d   }
 ||d |«      }|| j*                  |   d   |
<   |S c c}}
w c c}}
w )Nr   )rJ  rX  r   r-   zSpecialization data is for z but trying to preload for r  r  r  r  r
  r  rë   )rZ  rJ  rX  r  Útriton.languager¨   r   r[  rf  Úloadsr†  rQ   r'  rá   r%  r°   Úis_dtyper   r  r*   r%   rh  )r7   rG  rJ  rX  r  Útlr?  Údeserialized_objr  r  rë   rr   r  r  r  r  r
  r8   r  rv  s                       r:   ÚpreloadzJITFunction.preloadÅ  s³  € ß1ÛÝ$Ü—‘×1Ñ1Ó3ˆØŸ:™:Ð&9Ó:ÐØ˜FÑ# t§}¡}Ò4ÜØ-Ð.>¸vÑ.FÐ-GÐGbÐcg×cpÑcpÐbqÐrótð täœEÐ#3°OÑ#DÓEˆØ(¨Ñ9ˆô " -°Ó?÷
á��Uð  B§H¡H×$5Ñ$5°eÔ$<�—‘˜%”À%ÑGð
ˆ	ñ 
ô œÐ 0°Ñ >Ó?ˆ
Ø% lÑ3ˆ
Ü”S˜ ZÓ0Ó1ˆÜÐ)¨+Ñ6×<Ñ<Ó>Ó?ˆ	Ù˜˜i¨°EÓ:ˆð /¨yÑ9×?Ñ?ÓA÷
á��Uð ¤¨E´4Ô!8”�u”¸eÑCð
ˆñ 
ð ˜uÑ%ˆÙ˜˜d GÓ,ˆØ-3ˆ×Ñ˜6Ñ" 1Ñ% cÑ*Øˆùó!
ùó
s   Â9FÄ?(Fc                ó  — t        j                  | j                  «      }t        |t         j                  «      sJ ‚t        |j                  «      dk(  sJ ‚t        |j                  d   t         j                  «      sJ ‚|S )Nr�   r   )re   rŸ  r8   r*   ÚModuler"   ÚbodyÚFunctionDef)r7   Útrees     r:   rŸ  zJITFunction.parseå  s_   € Ü�y‰y˜Ÿ™Ó"ˆÜ˜$¤§
¡
Ô+Ð+Ð+Ü�4—9‘9‹~ Ò"Ð"Ð"Ü˜$Ÿ)™) A™,¬¯©Ô8Ð8Ð8Øˆr;   c                ó   — t        d«      ‚)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rQ   )r7   rx   rú   s      r:   Ú__call__zJITFunction.__call__ì  s   € ÜÐWÓXÐXr;   c                óV   •— |dk(  rt        d|› d�«      ‚t        t        | �  ||«       y )Nr8   zCannot set attribute 'zX' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)ÚAttributeErrorr+   rO   Ú__setattr__)r7   r-   rr   r9   s      €r:   r·  zJITFunction.__setattr__ï  s<   ø€ à�5Š=Ü Ð#9¸$¸ð @,ð "-ó .ð .ô 	Œk˜4Ñ,¨T°5Õ9r;   c                ó4   •— d| _         t        ‰| �	  d|«       y)zŽ
        The only method allowed to modify src.
        Bypasses the __setattr__ restriction by calling super().__setattr__ directly.
        Nr8   )r”  r+   r·  )r7   Únew_srcr9   s     €r:   r“  zJITFunction._unsafe_update_src÷  s   ø€ ð
 ˆŒ	Ü‰Ñ˜E 7Õ+r;   c                óP   — d| j                   › d| j                  j                  › d�S )NzJITFunction(ú:r  )rJ   r5  rœ   r?   s    r:   Ú__repr__zJITFunction.__repr__ÿ  s&   € Ø˜dŸk™k˜]¨!¨D¯G©G×,@Ñ,@Ð+AÀÐCÐCr;   )rš   zbool | None)NNNNNNN)rR   rB   rœ   r<  rQ  rU  r^  r  rI  r,   r�  rž   rT   rc   r  r­  rŸ  r´  r·  r“  r¼  rŸ   r    s   @r:   rO   rO   Û  s”   ø„ òð,
ð 
ó,
ò\(ò+òHòTFð mqØ;?ó<(ò|Lð ñ	ó ð	ð ñó ðò]òò@òYô:ô,öDr;   rO   c                 ó   — y r=   rÐ   r4  s    r:   Újitr¾    s   € àr;   ©rˆ  rI  ro  r¸   r¹   r`  rN   c                 ó   — y r=   rÐ   r¿  s          r:   r¾  r¾    s   € ð r;   c               ó@   ‡‡‡‡‡‡‡— dˆˆˆˆˆˆˆfd„}| � || «      S |S )a<  
    Decorator for JIT-compiling a function using the Triton compiler.

    :note: When a jit'd function is called, arguments are
        implicitly converted to pointers if they have a :code:`.data_ptr()` method
        and a `.dtype` attribute.

    :note: This function will be compiled and run on the GPU. It will only have access to:

           * python primitives,
           * builtins within the triton package,
           * arguments to this function,
           * other jit'd functions

    :param fn: the function to be jit-compiled
    :type fn: Callable
    c           
     ó¤   •— t        | «      sJ ‚t        j                  j                  rddlm}  || ‰‰‰‰‰‰‰¬«      S t        | ‰‰‰‰‰‰‰¬«      S )Nr�   )ÚInterpretedFunction)rˆ  r¸   r¹   r`  rN   rI  ro  )rS  r   re  Ú	interpretÚinterpreterrÃ  rO   )	r5  rÃ  r`  r¸   r¹   ro  rN   rI  rˆ  s	     €€€€€€€r:   Ú	decoratorzjit.<locals>.decorator8  sj   ø€ Ü˜Œ|Ðˆ|Ü�=‰=×"Ò"Ý8Ù& r°7ÐN_ØFdÐlqØ08¸tÐUdôfð fô ØØØ"3Ø/MØØ!ØØ /ô	ð 	r;   ©r5  r   rš   zJITFunction[T]rÐ   )	r5  rˆ  rI  ro  r¸   r¹   r`  rN   rÆ  s	    ``````` r:   r¾  r¾    s&   þ€ ÷:ó ð& 
€~Ù˜‹}Ðð Ðr;   c                  óF   — e Zd ZdZed„ «       Zd„ Zed„ «       Zed„ «       Zy)r¥  zr
    Can be used in place of real tensors when calling:
        kernel.warmup(MockTensor(torch.float32), ...)
    c                ól   — | j                   j                  dk(  r| j                  dk(  rt        | «      S | S )Nr°   Útorch)r9   rR   rB   r¥  )ry   s    r:   r¦  zMockTensor.wrap_dtype]  s.   € à�=‰=×!Ñ! WÒ,°·±À7Ò1JÜ˜c“?Ð"Øˆ
r;   c                ó   — || _         y r=   )r°   )r7   r°   s     r:   r,   zMockTensor.__init__c  s	   € Øˆ�
r;   c                  ó   — y©Nr   rÐ   rÐ   r;   r:   rÚ   zMockTensor.data_ptrf  ó   € àr;   c                  ó   — yrÍ  rÐ   rÐ   r;   r:   Ú	ptr_rangezMockTensor.ptr_rangej  rÎ  r;   N)	rR   rB   rœ   r�   Ústaticmethodr¦  r,   rÚ   rÐ  rÐ   r;   r:   r¥  r¥  W  sH   „ ñð
 ñó ðò
ð ñó ðð ñó ñr;   r¥  c                  óJ   — e Zd Zd„ Zd„ Zd„ Zdd„Zd„ Zd„ Zd„ Z	d„ Z
d	„ Zd
„ Zy)ÚTensorWrapperc                óš   — || _         || _        |j                  | _        |j                  | _        | j                  j                  | _        y r=   )r°   rç   Údatar?  Úshape)r7   rç   r°   s      r:   r,   zTensorWrapper.__init__q  s5   € ØˆŒ
ØˆŒ	Ø—I‘IˆŒ	Ø—k‘kˆŒØ—Y‘Y—_‘_ˆ�
r;   c                ó6   — | j                   j                  «       S r=   )rç   rÚ   r?   s    r:   rÚ   zTensorWrapper.data_ptrx  s   € Ø�y‰y×!Ñ!Ó#Ð#r;   c                ó4   —  | j                   j                  |Ž S r=   )rç   Ústride)r7   rx   s     r:   rÙ  zTensorWrapper.stride{  s   € Øˆt�y‰y×Ñ Ð&Ð&r;   c                ó<   — d| j                   › d| j                  › d�S )NzTensorWrapper[r>  r  )r°   rç   r?   s    r:   Ú__str__zTensorWrapper.__str__~  s   € Ø §
¡
˜|¨2¨d¯i©i¨[¸Ð:Ð:r;   c                ó6   — | j                   j                  «       S r=   )rç   Úelement_sizer?   s    r:   rÝ  zTensorWrapper.element_size�  s   € Ø�y‰y×%Ñ%Ó'Ð'r;   c                ó^   — t        | j                  j                  «       | j                  «      S r=   )rÓ  rç   Úcpur°   r?   s    r:   rß  zTensorWrapper.cpu„  s   € Ü˜TŸY™YŸ]™]›_¨d¯j©jÓ9Ð9r;   c                óN   — | j                   j                  |j                   «       y r=   )rç   Úcopy_)r7   Úothers     r:   rá  zTensorWrapper.copy_‡  s   € Ø�	‰	�‰˜Ÿ
™
Õ#r;   c                ó^   — t        | j                  j                  «       | j                  «      S r=   )rÓ  rç   Úcloner°   r?   s    r:   rä  zTensorWrapper.cloneŠ  s   € Ü˜TŸY™YŸ_™_Ó.°·
±
Ó;Ð;r;   c                ó`   — t        | j                  j                  |«      | j                  «      S r=   )rÓ  rç   Útor°   )r7   r?  s     r:   ræ  zTensorWrapper.to�  s    € Ü˜TŸY™YŸ\™\¨&Ó1°4·:±:Ó>Ð>r;   c                ó`   — t        | j                  j                  |«      | j                  «      S r=   )rÓ  rç   Ú	new_emptyr°   )r7   Úsizess     r:   rè  zTensorWrapper.new_empty�  s"   € Ü˜TŸY™Y×0Ñ0°Ó7¸¿¹ÓDÐDr;   NrÏ   )rR   rB   rœ   r,   rÚ   rÙ  rÛ  rÝ  rß  rá  rä  ræ  rè  rÐ   r;   r:   rÓ  rÓ  o  s5   „ ò%ò$ò'ó;ò(ò:ò$ò<ò?óEr;   rÓ  c                óø   — t        | t        «      r;|| j                  j                  k(  r| j                  S t        | j                  |«      S t	        | d«      rt        | |«      S t        dt        | «      › d�«      ‚)NrÚ   zCannot reinterpret a r3  )r*   rÓ  rç   r°   rà   r‘   rc   )rÛ   r°   s     r:   Úreinterpretrë  ”  sk   € Ü�&œ-Ô(Ø�F—K‘K×%Ñ%Ò%à—;‘;Ðô ! §¡¨eÓ4Ð4Ü	�˜Ô	$ä˜V UÓ+Ð+äÐ/´°V³¨~¸QÐ?Ó@Ð@r;   c                ód  — | }t        |t        «      s|j                  }t        |t        «      sŒ|j                  j                  j                  }t        j                  |j                  «      \  }}t        |«      D ].  \  }}|j                  «       j                  d«      sŒ&||z  } ||fS  ||fS )Nzdef )
r*   rO   r5  Ú__code__Úco_filenamerD   r‰  r‹  rª   rG   )r5  Úbase_fnÚ	file_nameÚlinesÚ
begin_lineÚidxÚlines          r:   Úget_jit_fn_file_linerõ  £  s§   € Ø€GÜ˜¤+Ô.Ø—*‘*ˆô ˜¤+Õ.à—
‘
×#Ñ#×/Ñ/€IÜ×.Ñ.¨w¯z©zÓ:Ñ€Eˆ:ô ˜uÓ%ò ‰	ˆˆTØ�:‰:‹<×"Ñ" 6Õ*Ø˜#ÑˆJØØ�jÐ Ð ð	ð �jÐ Ð r;   rÏ   )FrÇ  )rI  úOptional[Callable]ro  rö  r¸   úOptional[Iterable[int | str]]r¹   r÷  r`  úOptional[bool]rN   rø  rš   zCallable[[T], JITFunction[T]]r=   )r5  zOptional[T]rI  rö  ro  rö  r¸   r÷  r¹   r÷  r`  rø  rN   rø  rš   z4Union[JITFunction[T], Callable[[T], JITFunction[T]]])?Ú
__future__r   r   re   ri   r.   rD   r�   r�  rŒ  Úcollectionsr   Údataclassesr   Ú	functoolsr   Útypingr   r	   r
   r   r   r   r   r   r   r   Útriton.tools.tensor_descriptorr   Útypesr   rC   r   Úruntime.driverr   Ú_utilsr   r   r   r   rR   r"   rH   r   ÚNodeVisitorr   r¬   r´   ræ   rû   r÷   rþ   r   r  r1  r6  r8  rO   r¾  r¥  rÓ  rë  rõ  rÐ   r;   r:   ú<module>r     s  ðß ,Û 
Û Û Û Û Û 	Û Ý #Ý !Ý %ß d× d× då ;Ý Ý Ý #ß eÓ eàÐ.™3˜~Ó.Ð.Ð/€áˆCƒL€ôH!˜Ÿ™ô H!ó`D÷4/>ñ />ðd €	ØÐ ò6ór@ô	Z�g˜a‘jô 	Zò	ò7*òt0ð ÷ð ó ðôeD�/ !Ñ$ô eDðZ	 
òó 
ðð 
ð Ø#Ø*.Ø7;ØDHØ Ø#ñ
ð ð
ð (ð	
ð
 5ð
ð %Bð
ð ð
ð ð
ð #ò
ó 
ð
ð ð4ð Ø#Ø*.Ø7;ØDHØ Ø#ñ4Øð4ð ð	4ð
 (ð4ð 5ð4ð %Bð4ð ð4ð ð4ð :ó4÷xñ ÷0"Eñ "EòJAó!r;   