§
    9˜ºhž�  ã            	      ó²  — 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j,        ¦  «        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=d0„¦   «         Z:	 d>dddddddd%œd?d3„Z: G d4„ d5¦  «        Z; G d6„ d7¦  «        Z<d8„ Z=d9„ Z>dS )@é    )Ú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                  ó|   ‡ — 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).
    ÚreturnÚNonec                óø   •— t          ¦   «                              ¦   «          || _        t          j        |                     d¦  «        ¦  «        | _        || _        || _        h d£| _	        i | _
        d| _        d S )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/   r4   r5   ÚsrcÚ	__class__s        €úZ/var/www/html/services/rag-service/venv/lib/python3.11/site-packages/triton/runtime/jit.pyr.   zDependenciesFinder.__init__)   s{   ø€ Ý‰Œ×ÒÑÔÐØˆŒ	Ý”n S§Z¢Z°Ñ%8Ô%8Ñ9Ô9ˆŒð ˆŒØ"ˆŒð*
ð *
ð *
ˆÔ&ð. TVˆÔà*/ˆÔ'Ð'Ð'ó    c                ó4   — | j                              ¦   «         S ©N)r3   Ú	hexdigest©r9   s    r<   ÚretzDependenciesFinder.retN   s   € àŒ{×$Ò$Ñ&Ô&Ð&r=   c                óŽ   — t          j        |j        ¦  «        rdS t          |dd¦  «        }|                     t
          ¦  «        S )NTÚ
__module__Ú )ÚinspectÚ	isbuiltinÚfuncr+   Ú
startswithÚTRITON_MODULE)r9   ÚnoderH   Úmodules       r<   Ú_is_triton_builtinz%DependenciesFinder._is_triton_builtinR   sA   € ÝÔ˜TœYÑ'Ô'ð 	Ø�4Ý˜˜|¨RÑ0Ô0ˆØ× Ò ¥Ñ/Ô/Ð/r=   c                ó2  — t          |t          ¦  «        �r | j                             ¦   «         |j                             ¦   «         z  D ]V}|\  }}| j        |         \  }}|j        |         \  }}||k    r)t	          d|› d|› d| j        › d|j        › d|› d�¦  «        ‚ŒW| j                             |j        ¦  «         |j        }|t          t          |dd¦  «        ¦  «        z  }| j                             |                     d	¦  «        ¦  «         d S d S )
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,   ÚJITFunctionr7   ÚkeysÚRuntimeErrorr/   Ú__name__ÚupdateÚ	cache_keyÚstrr+   r3   r2   )r9   rH   ÚkÚvar_nameÚ_Úv1Úv2Úfunc_keys           r<   Ú_update_hashzDependenciesFinder._update_hashX   se  € Ý�d�KÑ(Ô(ñ 	9ð Ô*×/Ò/Ñ1Ô1°DÔ4I×4NÒ4NÑ4PÔ4PÑPð ð �Ø‘�˜!ØÔ-¨aÔ0‘��AØÔ-¨aÔ0‘��AØ˜’8�8Ý&ð T¨8ð  Tð  TÀð  Tð  TÐTXÔT]ð  Tð  TÐrvÔrð  Tð  Tð  Y[ð  Tð  Tð  Tñô ð ð ð Ô!×(Ò(¨Ô)>Ñ?Ô?Ð?à”~ˆHØ��G D¨*°eÑ<Ô<Ñ=Ô=Ñ=ˆHØŒK×Ò˜xŸš¨wÑ7Ô7Ñ8Ô8Ð8Ð8Ð8ð	9ð 	9r=   c                óÚ  ‡ — t          |j        ¦  «        t          j        u r|j        S |j        ‰ j        v rd S ˆ fd„} ||j        ¦  «        \  }}|�ƒ‰ j        s|t          |¦  «        t          urft          |t          ¦  «        sQt          |dd¦  «        s@|j        ‰ j        vr2t          j        |¦  «        |f‰ j        |j        t	          |¦  «        f<   ‰                      |¦  «         |S )Nc                ó    •— ‰j                              | d ¦  «        }|�	|‰j         fS ‰j                             | d ¦  «        }|�	|‰j        fS dS )N)NN)r4   Úgetr5   )r/   Úvalr9   s     €r<   Úname_lookupz2DependenciesFinder.visit_Name.<locals>.name_lookupr   sZ   ø€ Ø”,×"Ò" 4¨Ñ.Ô.ˆCØˆØ˜DœLÐ(Ð(Ø”.×$Ò$ T¨4Ñ0Ô0ˆCØˆØ˜DœNÐ*Ð*Ø�:r=   Ú__triton_builtin__F)ÚtypeÚctxÚastÚStoreÚidÚlocal_namesr8   r   r,   rQ   r+   r6   Úcopyr7   r^   )r9   rK   rc   rb   Úvar_dicts   `    r<   Ú
visit_NamezDependenciesFinder.visit_Namej   s
  ø€ Ý�”‰>Œ>�SœYÐ&Ð&Ø”7ˆNàŒ7�dÔ&Ð&Ð&à�4ð	ð 	ð 	ð 	ð 	ð $˜ D¤GÑ,Ô,‰ˆˆXð
 ˆOð Ô7ð	 õ ˜‘I”I¥ZÐ/Ð/õ # 3­Ñ4Ô4ð 0õ >EÀSÐJ^Ð`eÑ=fÔ=fð 0ð ”G 4Ô#AÐAÐAÝ>B¼iÈ¹n¼nÈhÐ=WˆDÔ! 4¤7­B¨x©L¬LÐ"9Ñ:à×Ò˜#ÑÔÐØˆ
r=   c                ó*   ‡ — ˆ fd„|j         D ¦   «         S )Nc                ó:   •— g | ]}‰                      |¦  «        ‘ŒS © )Úvisit)Ú.0Úeltr9   s     €r<   ú
<listcomp>z2DependenciesFinder.visit_Tuple.<locals>.<listcomp>”   s#   ø€ Ð5Ð5Ð5 C�—
’
˜3‘”Ð5Ð5Ð5r=   )Úelts©r9   rK   s   ` r<   Úvisit_TuplezDependenciesFinder.visit_Tuple‘   s!   ø€ ð 6Ð5Ð5Ð5¨4¬9Ð5Ñ5Ô5Ð5r=   c                óf  — |                       |j        ¦  «        }t          |t          j        ¦  «        r4|                       |j        ¦  «        }t          |t          j        ¦  «        °4|�t          |dd¦  «        t          k    rd S t          ||j        ¦  «        }|                      |¦  «         |S )NrT   rE   )	rq   Úvaluer,   rg   Ú	Attributer+   rJ   Úattrr^   )r9   rK   ÚlhsrB   s       r<   Úvisit_Attributez"DependenciesFinder.visit_Attribute–   sœ   € Ø�jŠj˜œÑ$Ô$ˆÝ˜�cœmÑ,Ô,ð 	(Ø—*’*˜SœYÑ'Ô'ˆCõ ˜�cœmÑ,Ô,ð 	(àˆ;�7 3¨
°BÑ7Ô7½=ÒHÐHØ�4Ý�c˜4œ9Ñ%Ô%ˆØ×Ò˜#ÑÔÐØˆ
r=   c                óf   — d„ |j         j         D ¦   «         | _        |                      |¦  «         d S )Nc                ó   — h | ]	}|j         ’Œ
S rp   ©Úarg)rr   r�   s     r<   ú	<setcomp>z7DependenciesFinder.visit_FunctionDef.<locals>.<setcomp>¢   s   € Ð>Ð>Ð>¨˜CœGÐ>Ð>Ð>r=   )Úargsrj   Úgeneric_visitrv   s     r<   Úvisit_FunctionDefz$DependenciesFinder.visit_FunctionDef    s6   € à>Ð>¨t¬y¬~Ð>Ñ>Ô>ˆÔØ×Ò˜4Ñ Ô Ð Ð Ð r=   c                ó.  ‡ — ˆ fd„}t          j        |j        |j        |j        r|j        gng |j        ¦  «        D ]}‰                      |¦  «         Œ ||j        ¦  «         |j        �‰                      |j        ¦  «          ||j	        ¦  «         d S )Nc                óŠ   •— 	 ‰j         rJ ‚d‰_         | D ]}|�‰                     |¦  «         Œ	 d‰_         d S # d‰_         w xY w)NTF)r8   rq   )ÚdefaultsÚexprr9   s     €r<   Úvisit_defaultsz:DependenciesFinder.visit_arguments.<locals>.visit_defaults±   sp   ø€ ð8ØÔ:Ð:Ð:Ð:Ø26�Ô/Ø$ð )ð )�DØÐ'ØŸ
š
 4Ñ(Ô(Ð(øð)ð 38�Ô/Ð/Ð/ø°%�Ô/Ð7Ð7Ð7Ð7s	   ƒ,9 ¹	A)
Ú	itertoolsÚchainÚposonlyargsrƒ   ÚvarargÚ
kwonlyargsrq   Úkw_defaultsÚkwargrˆ   )r9   rK   rŠ   r�   s   `   r<   Úvisit_argumentsz"DependenciesFinder.visit_arguments¥   s°   ø€ ð	8ð 	8ð 	8ð 	8ð 	8õ ”? 4Ô#3°T´YÐQUÔQ\Ð@dÀÄÀÀÐbdÐfjÔfuÑvÔvð 	ð 	ˆCØ�JŠJ�s‰OŒOˆOˆOàˆ�tÔ'Ñ(Ô(Ð(àŒ:Ð!Ø�JŠJ�t”zÑ"Ô"Ð"àˆ�t”}Ñ%Ô%Ð%Ð%Ð%r=   c                óÌ   — |                       |¦  «        }t          |t          ¦  «        r| xj        t	          |¦  «        z  c_        d S | j                             |¦  «         d S r?   )rq   r,   r'   rj   ÚsetÚadd)r9   rK   Útargets      r<   ÚvisitAssnTargetz"DependenciesFinder.visitAssnTargetÅ   sd   € ð —’˜DÑ!Ô!ˆÝ�f�dÑ#Ô#ð 	)ØÐÔ¥ F¡¤Ñ+ÐÔÐÐàÔ× Ò  Ñ(Ô(Ð(Ð(Ð(r=   c                ó¾   — t          |j        ¦  «        dk    rt          d¦  «        ‚|                      |j        d         ¦  «         |                      |¦  «         d S )Né   z2Simultaneous multiple assignment is not supported.r   )r$   ÚtargetsÚ	TypeErrorr—   r„   rv   s     r<   Úvisit_AssignzDependenciesFinder.visit_AssignÎ   s^   € ÝˆtŒ|ÑÔ Ò!Ð!õ
 ÐPÑQÔQÐQà×Ò˜Tœ\¨!œ_Ñ-Ô-Ð-ð 	×Ò˜4Ñ Ô Ð Ð Ð r=   c                ód   — |                       |j        ¦  «         |                      |¦  «         d S r?   ©r—   r–   r„   rv   s     r<   Úvisit_AnnAssignz"DependenciesFinder.visit_AnnAssignÛ   ó4   € Ø×Ò˜Tœ[Ñ)Ô)Ð)ð 	×Ò˜4Ñ Ô Ð Ð Ð r=   c                ód   — |                       |j        ¦  «         |                      |¦  «         d S r?   rž   rv   s     r<   Ú	visit_ForzDependenciesFinder.visit_Forá   r    r=   )r   r    )rT   rD   Ú__qualname__Ú__doc__r.   ÚpropertyrB   rM   r^   rm   rw   r}   r…   r’   r—   rœ   rŸ   r¢   Ú__classcell__©r;   s   @r<   r   r      s
  ø€ € € € € ð	ð 	ð#0ð #0ð #0ð #0ð #0ð #0ðJ ð'ð 'ñ „Xð'ð0ð 0ð 0ð9ð 9ð 9ð$%ð %ð %ðN6ð 6ð 6ð
ð ð ð!ð !ð !ð
&ð &ð &ð@)ð )ð )ð!ð !ð !ð!ð !ð !ð!ð !ð !ð !ð !ð !ð !r=   r   r   rW   c                óŠ  — dd l mc m} t          | t          ¦  «        �r|                      ¦   «         } |                      d¦  «        rH|                      d¦  «        } t          | ¦  «        } |                      d¦  «        sJ ‚d| dd …         z   S |  	                    d¦  «        rdt          | d d…         ¦  «        z   S |                      d¦  «        rdt          | dd …         ¦  «        z   S |                      d¦  «        r"t          |                      d¦  «        ¦  «        S nut          | |j
        ¦  «        rdt          | j        ¦  «        › �S t          | |j        ¦  «        r| j        } n,t          | t          ¦  «        r| j        } nt	          | ¦  «        } t!          j        |                      d	d
¦  «        | ¦  «        S )Nr   zconst ÚconstÚ*ú*kr™   éÿÿÿÿztl.Ú_trE   )Útriton.language.coreÚlanguageÚcorer,   rW   ÚstriprI   ÚremoveprefixÚ_normalize_tyÚendswithÚpointer_typeÚ
element_tyÚdtyper/   re   rT   r   ra   Úreplace)Útyr°   s     r<   r³   r³   í   sº  € Ø'Ð'Ð'Ð'Ð'Ð'Ð'Ð'Ð'Ý�"•cÑÔñ Ø�XŠX‰ZŒZˆØ�=Š=˜Ñ"Ô"ð 	!Ø—’ Ñ)Ô)ˆBÝ˜rÑ"Ô"ˆBØ—=’= Ñ%Ô%Ð%Ð%Ð%Ø˜"˜Q˜R˜Rœ&‘=Ð Ø�;Š;�sÑÔð 	0Ø� r¨#¨2¨#¤wÑ/Ô/Ñ/Ð/Ø�=Š=˜ÑÔð 	/Ø� r¨!¨"¨"¤vÑ.Ô.Ñ.Ð.Ø�=Š=˜ÑÔð 	9Ý  §¢°Ñ!7Ô!7Ñ8Ô8Ð8ð	9å	�B˜Ô)Ñ	*Ô	*ð Ø1•= ¤Ñ/Ô/Ð1Ð1Ð1Ý	�B˜œ
Ñ	#Ô	#ð ØŒWˆˆÝ	�B�Ñ	Ô	ð ØŒ[ˆˆå�‰WŒWˆÝ%Ô)¨"¯*ª*°T¸2Ñ*>Ô*>ÀÑCÔ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dS )ÚKernelParamzBRepresents a parameter (name plus metadata) to a @jit'ed function.Únumr#   Úparamúinspect.ParameterÚdo_not_specializeÚboolÚdo_not_specialize_on_alignmentc                ó>   — || _         || _        || _        || _        d S r?   )r¼   Ú_paramr¿   rÁ   )r9   r¼   r½   r¿   rÁ   s        r<   r.   zKernelParam.__init__
  s&   € àˆŒØˆŒØ!2ˆÔØ.LˆÔ+Ð+Ð+r=   c                ó   — | j         j        S r?   )rÃ   r/   rA   s    r<   r/   zKernelParam.name  s   € àŒ{ÔÐr=   r   rW   c                óŽ   — | j         j        r| j         j        t          j        j        k    rdS t          | j         j        ¦  «        S )NrE   )rÃ   Ú
annotationrF   Ú	ParameterÚemptyr³   rA   s    r<   rÆ   zKernelParam.annotation  s<   € àŒ{Ô%ð 	¨¬Ô)?Å7ÔCTÔCZÒ)ZÐ)ZØ�2Ý˜Tœ[Ô3Ñ4Ô4Ð4r=   c                óä   — | j         }|                     d¦  «        r|dd …         }n|                     d¦  «        r
|dd …         }|t          t          j        ¦   «         ¦  «        v r| j         S dS )Nr«   r   rª   r™   rE   )rÆ   rI   r”   r   Úvalues)r9   Úas     r<   Úannotation_typezKernelParam.annotation_type  sw   € àŒOˆØ�<Š<˜ÑÔð 	Ø�!�"�"”ˆAˆAØ�\Š\˜#ÑÔð 	Ø�!�"�"”ˆAØ•Õ.Ô5Ñ7Ô7Ñ8Ô8Ð8Ð8Ø”?Ð"Øˆrr=   c                ó   — d| j         v S ©NÚ	constexpr)rÆ   rA   s    r<   Úis_constexprzKernelParam.is_constexpr&  s   € à˜dœoÐ-Ð-r=   c                óZ   — | j         rdS d| j        v p| j                             d¦  «        S )NFr©   r«   )rÐ   rÆ   rI   rA   s    r<   Úis_constzKernelParam.is_const*  s4   € àÔð 	Ø�5Ø˜$œ/Ð)ÐM¨T¬_×-GÒ-GÈÑ-MÔ-MÐMr=   c                ó   — | j         j        S r?   )rÃ   ÚdefaultrA   s    r<   rÔ   zKernelParam.default0  s   € àŒ{Ô"Ð"r=   c                ó@   — | j         j        t          j        j        k    S r?   )rÃ   rÔ   rF   rÇ   rÈ   rA   s    r<   Úhas_defaultzKernelParam.has_default4  s   € àŒ{Ô"¥gÔ&7Ô&=Ò=Ð=r=   N)r¼   r#   r½   r¾   r¿   rÀ   rÁ   rÀ   ©r   rW   )rT   rD   r£   r¤   r.   r   r/   rÆ   rÌ   rÐ   rÒ   r¥   rÔ   rÖ   rp   r=   r<   r»   r»     sü   € € € € € ØLÐLðMð Mð Mð Mð ð ð  ñ „_ð ð ð5ð 5ð 5ñ „_ð5ð
 ðð ð ñ „_ðð ð.ð .ñ „_ð.ð ðNð Nñ „_ðNð
 ð#ð #ñ „Xð#ð ð>ð >ñ „Xð>ð >ð >r=   r»   c                ó8   ‡ ‡‡‡— ddl mŠ ddlmŠ dˆˆˆ ˆfd„	Š‰S )	Nr   ©rÏ   r   r   FTc                ó   •‡ — ‰ €dS t          ‰ t          ¦  «        rdS t          ‰ t          ¦  «        r@|r ‰‰ d|¬¦  «        nd }‰ dk    r|rdS d‰ k    r
‰ dk    rd	|fS d
‰ k    r
‰ dk    rd|fS d|fS t          ‰ t          ¦  «        rdS t	          ‰ d¦  «        rh‰ j        |f}t                               |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dS t          ‰ t          ¦  «        rAˆfd„‰ D ¦   «         }ˆ fd„} |d„ |D ¦   «         ¦  «        }	 |d„ |D ¦   «         ¦  «        }
|	|
fS t          ‰ t          ¦  «        rLt	          ‰ j        d¦  «        sJ ‚t          ‰ j        j        ¦  «        }d|› t          ‰ j        ¦  «        › d�d fS t          ‰ ‰¦  «        rTt	          ‰ j        d¦  «        sJ ‚t          ‰ j        j        ¦  «        }d|› t          ‰ j        ¦  «        › d‰ j        ›d�d fS t#          dt%          ‰ ¦  «        z  ¦  «        ‚)N)rÏ   N)Úu1Nr#   )Úalignr™   )rÏ   r™   i   €iÿÿÿÚi32l            l   ÿÿÿÿ Úu64Úi64)Úfp32NÚdata_ptrr«   rª   r   ÚtensorrÏ   Útma_desc_cpu_ptr)Ú	nvTmaDescNc                ó&   •— g | ]} ‰|¦  «        ‘ŒS rp   rp   )rr   ÚxÚspecialize_impls     €r<   rt   zCcreate_specialize_impl.<locals>.specialize_impl.<locals>.<listcomp>c  s#   ø€ Ð4Ð4Ð4¨1�O�O AÑ&Ô&Ð4Ð4Ð4r=   c                óf   •— t          ‰d¦  «        r t          ‰¦  «        | Ž nt          | ¦  «        S )NÚ_fields)Úhasattrre   Útuple)Úvalsr�   s    €r<   ú<lambda>zAcreate_specialize_impl.<locals>.specialize_impl.<locals>.<lambda>d  s2   ø€ ½'À#ÀyÑ:QÔ:QÐ&b i¥d¨3¡i¤i°Ð&6Ð&6ÕW\Ð]aÑWbÔWb€ r=   c                ó   — g | ]
}|d          ‘ŒS ©r   rp   ©rr   ræ   s     r<   rt   zCcreate_specialize_impl.<locals>.specialize_impl.<locals>.<listcomp>e  s   € Ð1Ð1Ð1 q˜a œdÐ1Ð1Ð1r=   c                ó   — g | ]
}|d          ‘ŒS ©r™   rp   rð   s     r<   rt   zCcreate_specialize_impl.<locals>.specialize_impl.<locals>.<listcomp>f  s   € Ð2Ð2Ð2¨˜q œtÐ2Ð2Ð2r=   ztensordesc<ú>ú,zUnsupported type: %s)r,   rÀ   r#   r(   rê   r·   Ú	dtype2strra   r   rQ   rV   rë   r   Úbaser'   Úblock_shapeÚlayoutr›   re   )r�   rÒ   Úspecialize_valuerÜ   ÚkeyÚdskÚresÚspecÚ
make_tupleÚtysrR   ÚinnerÚGluonTensorDescriptorrÏ   Úspecialize_extrarç   s   `           €€€€r<   rç   z/create_specialize_impl.<locals>.specialize_implB  s  øø€ Øˆ;Ø&Ð&Ý˜�TÑ"Ô"ð ,	@Ø�<Ý˜�SÑ!Ô!ð *	@Ø?OÐYÐ"Ð" 3¨°UÐ;Ñ;Ô;Ð;ÐUYˆCØ�aŠxˆxÐ,ˆxØ'Ð'Ø˜S’� S¨IÒ%5Ð%5Ø˜s�|Ð#Ø˜#’� #¨Ò"2Ð"2Ø˜s�|Ð#à˜s�|Ð#Ý˜�UÑ#Ô#ð  	@Ø!�>Ý�S˜*Ñ%Ô%ð 	@à”9˜hÐ'ˆCÝ—-’-  TÑ*Ô*ˆCØˆ{Ø" 1œvÐ.�t�t¨3Õ2DÀSÈÄVÑ2LÔ2LÑL�Ø!$•	˜#‘ØBRÐ\Ð"Ð" 3¨¸Ð>Ñ>Ô>Ð>ÐX\ˆCØ˜�:ÐÝ˜�[Ñ)Ô)ð 	@Ø ¤Ð/Ð/Ý˜˜YÑ'Ô'ð 	@Ø Ð%Ð%Ý�SÐ,Ñ-Ô-ð 	@Ø&Ð&Ý˜�UÑ#Ô#ð 	@Ø4Ð4Ð4Ð4°Ð4Ñ4Ô4ˆDØbÐbÐbÐbˆJØ�*Ð1Ð1¨DÐ1Ñ1Ô1Ñ2Ô2ˆCØ�:Ð2Ð2¨TÐ2Ñ2Ô2Ñ3Ô3ˆDØ˜�;ÐÝ˜Õ-Ñ.Ô.ð 		@Ý˜3œ8 ZÑ0Ô0Ð0Ð0Ð0Ý& s¤x¤~Ñ6Ô6ˆEØA %ÐA­¨c¬oÑ)>Ô)>ÐAÐAÐAÀ4ÐHÐHÝ˜Ð2Ñ3Ô3ð 	@Ý˜3œ8 ZÑ0Ô0Ð0Ð0Ð0Ý& s¤x¤~Ñ6Ô6ˆEØP %ÐP­¨c¬oÑ)>Ô)>ÐPÐPÀÄÐPÐPÐPÐRVÐWÐWåÐ2µT¸#±Y´YÑ>Ñ?Ô?Ð?r=   )FTT)r¯   rÏ   Ú'triton.experimental.gluon.nvidia.hopperr   )r  r  rÏ   rç   s   `@@@r<   Úcreate_specialize_implr  =  si   øøøø€ à$Ð$Ð$Ð$Ð$Ð$ØaÐaÐaÐaÐaÐað/@ð /@ð /@ð /@ð /@ð /@ð /@ð /@ð /@ðb Ðr=   Fc                óÂ   — t          t          ¦  «        dk    r(t                               t          d„ ¦  «        ¦  «         t          d         } || |¬¦  «        d         S )Nr   c                ó   — d S r?   rp   )rZ   Úkwargss     r<   rí   zmangle_type.<locals>.<lambda>x  s   € ÐPT€ r=   )rù   )r$   Úspecialize_impl_cacheÚappendr  )r�   Ú
specializerç   s      r<   Úmangle_typer  v  s[   € Ý
Õ Ñ!Ô! QÒ&Ð&Ý×$Ò$Õ%;Ð<TÐ<TÑ%UÔ%UÑVÔVÐVÝ+¨AÔ.€OØˆ?˜3°Ð<Ñ<Ô<¸QÔ?Ð?r=   c                  ó"   — e Zd ZU ded<   dd„ZdS )ÚKernelInterfacer   Úrunr   c                ó   ‡ ‡— ˆˆ 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  )rƒ   r  r  r9   s     €€r<   rí   z-KernelInterface.__getitem__.<locals>.<lambda>†  s    ø€  x t¤xÈ$°TÀ%Ð'YÐ'YÐRXÐ'YÐ'Y€ r=   rp   )r9   r  s   ``r<   Ú__getitem__zKernelInterface.__getitem__€  s   øø€ ð ZÐYÐYÐYÐYÐYr=   N)r   r   )rT   rD   r£   Ú__annotations__r  rp   r=   r<   r  r  }  s9   € € € € € € Ø
€F€F�FðZð Zð Zð Zð Zð Zr=   r  c           	     ó€  — d„ |                      ¦   «         D ¦   «         }dd l}| |d„ |                     ¦   «         D ¦   «         t          |                     ¦   «         ¦  «        d„ |                     ¦   «         D ¦   «         t          |                     ¦   «         ¦  «        |j        |dœ}|                     |¦  «        }|S )Nc                óX   — i | ]'\  }}||j         j        d k    rt          |¦  «        n|“Œ(S ©r·   )r;   rT   rW   ©rr   rú   ry   s      r<   ú
<dictcomp>z1serialize_specialization_data.<locals>.<dictcomp>‹  s:   € ÐwÐwÐwÑWaÐWZÐ\a� E¤OÔ$<ÀÒ$GÐ$G•c˜%‘j”j�jÈUÐwÐwÐwr=   r   c                ó,   — g | ]}t          |¦  «        ‘ŒS rp   ©r'   rð   s     r<   rt   z1serialize_specialization_data.<locals>.<listcomp>Ž  s   € Ð?bÐ?bÐ?bÈAÅÀQÁÄÐ?bÐ?bÐ?br=   c                ó,   — g | ]}t          |¦  «        ‘ŒS rp   r  rð   s     r<   rt   z1serialize_specialization_data.<locals>.<listcomp>�  s   € Ð0OÐ0OÐ0O¸Qµ°a±´Ð0OÐ0OÐ0Or=   )r/   Ú	signatureÚconstant_keysÚconstant_valsÚ
attrs_keysÚ
attrs_valsÚoptionsrú   )ÚitemsÚjsonrR   r'   rÊ   Ú__dict__Údumps)	r/   r  Ú	constantsÚattrsr#  rú   r%  ÚobjÚserialized_objs	            r<   Úserialize_specialization_datar,  Š  sÁ   € ØwÐwÐen×etÒetÑevÔevÐwÑwÔw€IØ€K€K€Kà 9Ð?bÐ?bÐQZ×Q_ÒQ_ÑQaÔQaÐ?bÑ?bÔ?bÝˆY×ÒÑÔÑ Ô Ð0OÐ0OÀ%Ç*Â*Á,Ä,Ð0OÑ0OÔ0OÕ_cÐdi×dpÒdpÑdrÔdrÑ_sÔ_sØÔ#¨Cðð €Cð
 —Z’Z ‘_”_€NØÐr=   c           
     óv  — t          | j        ¦  «        t          |¦  «        k    sJ ‚g }t          | j                             ¦   «         |¦  «        D ]î\  }}|j        r|                     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"|                     d|j	        › d|	› d�¦  «         Œ¹|                     d|j	        › d�¦  «         ŒØ|                     |	› ¦  «         Œïd„ }
dd                     t          t          |
| j                             ¦   «         ¦  «        ¦  «        dgz   ¦  «        › dd                     d„ | j                             ¦   «         D ¦   «         ¦  «        › dd                     |¦  «        › d�}d„ | j                             ¦   «         D ¦   «         }t           |d<   t#          |j        ¦  «        |d<   t'          ||¦  «         |d         S )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                ót   — | d         j         t          j        j        u r| d         n| d         › d| d         › �S )Nr™   r   z	=default_©rÔ   rF   rÇ   rÈ   )ræ   s    r<   rí   z0create_function_from_signature.<locals>.<lambda>·  s=   € ˜A˜aœDœL­GÔ,=Ô,CÐCÐC�A�a”D�DÈAÈaÌDÐIaÐIaÐ[\Ð]^Ô[_ÐIaÐIa€ r=   z
def dynamic_func(z	**optionsz):
    params = {c                ó   — g | ]
}d |› d|› �‘ŒS )ú'z': rp   )rr   r/   s     r<   rt   z2create_function_from_signature.<locals>.<listcomp>»  s)   € ÐQÐQÐQ°4Ð.˜tÐ.Ð.¨Ð.Ð.ÐQÐQÐQr=   z}
    specialization = [rô   z-]
    return params, specialization, options
c                óZ   — i | ](\  }}|j         t          j        j        u¯d |› �|j         “Œ)S )Údefault_r5  )rr   r/   r½   s      r<   r  z2create_function_from_signature.<locals>.<dictcomp>À  sE   € ð ð ð áˆD�%ØŒ=¥Ô 1Ô 7Ð7Ð7ð 	�4ÐÐ˜5œ=à7Ð7Ð7r=   rQ   rç   Údynamic_func)r$   Ú
parametersÚziprR   rÐ   r	  rÒ   r¿   rÁ   rÌ   r,   rW   Újoinr'   Úmapr$  rQ   r  Úget_arg_specializationÚexec)ÚsigÚkparamsÚbackendÚspecializationr/   ÚkprÒ   r
  rÜ   rB   r�   Ú	func_bodyÚfunc_namespaces                r<   Úcreate_function_from_signaturerH  –  s²  € õ ˆsŒ~ÑÔ¥# g¡,¤,Ò.Ð.Ð.Ð.à€Nå˜œ×+Ò+Ñ-Ô-¨wÑ7Ô7ð 0ð 0‰ˆˆbØŒ?ð 	0Ø×!Ò!Ð":°4Ð":Ð":Ð":Ñ;Ô;Ð;Ð;à!#¤Ð9�v�v°'ˆHØ$&Ô$8ÐD˜˜¸fˆJØ!Ô@ÐL�G�GÀfˆEØO TÐOÐO¨XÐOÐO¸ÐOÐOÀuÐOÐOÐOˆCØÔ!ð 0Ý˜bÔ0µ#Ñ6Ô6ð +ØÔ)¨TÒ1Ð1°RÔ5GÈÈÈÔ5KÈ|Ð5[Ð5[à%*˜
Øð MØ"×)Ò)Ð*R¨rÔ/AÐ*RÐ*RÈÐ*RÐ*RÐ*RÑSÔSÐSÐSð #×)Ò)Ð*K¨rÔ/AÐ*KÐ*KÐ*KÑLÔLÐLÐLà×%Ò%¨ hÑ/Ô/Ð/Ð/ð bÐ
a€CðØ—)’)�D¥ S¨#¬.×*>Ò*>Ñ*@Ô*@Ñ!AÔ!AÑBÔBÀkÀ]ÑRÑSÔSðð à—	’	ÐQÐQ¸3¼>×;NÒ;NÑ;PÔ;PÐQÑQÔQÑRÔRðð ð —x’x Ñ/Ô/ðð ð €Iðð àœ>×/Ò/Ñ1Ô1ðñ ô €Nõ %0€N�=Ñ!Ý(>¸wÔ?]Ñ(^Ô(^€NÐ$Ñ%õ 	ˆ�NÑ#Ô#Ð#ð ˜.Ô)Ð)r=   c                ó$   — | j         › d| j        › �S )Nú.)rD   r£   ©Úfns    r<   Úget_full_namerM  Ð  s   € ØŒmÐ/Ð/˜bœoÐ/Ð/Ð/r=   c                  ó.   — e Zd ZU ded<   ded<   ded<   dS )ÚJitFunctionInfor   rL   rW   r/   rQ   Újit_functionN)rT   rD   r£   r  rp   r=   r<   rO  rO  Ô  s3   € € € € € € àÐÐÑØ€I€I�IØÐÐÑÐÐr=   rO  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 )rQ   c                ó   — dS )NFrp   rA   s    r<   Úis_gluonzJITFunction.is_gluonÝ  s   € Øˆur=   r   úbool | Nonec	                ó  — |sd S | j         j        }	| j         j        }
d                     d„ t	          | j        |d         ¦  «        D ¦   «         ¦  «        }|	› 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 )Nr1  c                ó,   — g | ]\  }}|j         › d |› �‘ŒS )z: ©r/   )rr   r½   r¹   s      r<   rt   z*JITFunction._call_hook.<locals>.<listcomp>ð  s,   € Ð_Ð_Ð_¹¸À %¤*Ð4Ð4°Ð4Ð4Ð_Ð_Ð_r=   r™   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ú   ÚreprrL  ÚcompileÚis_manual_warmupÚalready_compiled)rL  r£   rD   r=  r<  ÚparamsrZ  r[  r\  r]  r^  rM  r,  r_  rO  )r9   Úhookrú   r  rY  r(  r#  r`  rb  r/   rL   Ú	arg_reprsrc  Ú	full_namera  r  s                   r<   Ú
_call_hookzJITFunction._call_hookà  s¢  € ð ð 	Ø�4àŒwÔ#ˆØ”Ô#ˆØ—I’IÐ_Ð_ÅcÈ$Ì+ÐWZÐ[\ÔW]ÑF^ÔF^Ð_Ñ_Ô_Ñ`Ô`ˆ	Øð  kð  k 7Ô#4ð  kð  kÀÔAQð  kð  kÐ`gÔ`rð  kð  kð  HOô  H`ð  kð  kð  |Cô  |[ð  kð  kð  _hð  kð  kð  kˆÝ! $¤'Ñ*Ô*ˆ	å;¸IÀyÐR[Ð]dÐefÔ]gÐipÐruÑvÔvÐð #ØØ"Ø Ô*ØÔ(Ø!Ô,Ø 'Ô 8Ø'.Ô'FØ"Ô.ØØ#6Ø"ð
ð 
ˆð ˆtØØÝ˜v t¨TÑ2Ô2Ø˜CÐ* 6Ð*Ø&Ø"ð
ñ 
ô 
ð 	
r=   c                ó\   — t          |¦  «        sJ ‚| j                             |¦  «         dS )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	  )r9   rh  s     r<   Úadd_pre_run_hookzJITFunction.add_pre_run_hook  s3   € õ
 ˜‰~Œ~ÐÐÐØÔ×!Ò! $Ñ'Ô'Ð'Ð'Ð'r=   c                óØ   — ddl m}m}m}m} t
          j                             ¦   «         } ||¦  «        }|| _        || _        || _        t          | j	        | j
        |¦  «        }i |||fS )z1
        Precompute as much as possible.
        r   )ÚCompiledKernelrd  Ú	ASTSourceÚmake_backend)Úcompilerrq  rd  rr  rs  r   ÚactiveÚget_current_targetrH  r  rg  )r9   rq  rd  rr  rs  r–   rC  Úbinders           r<   Úcreate_binderzJITFunction.create_binder  s…   € ð 	PÐOÐOÐOÐOÐOÐOÐOÐOÐOÐOÐOÝ”×1Ò1Ñ3Ô3ˆØ�,˜vÑ&Ô&ˆØ,ˆÔØˆŒØ"ˆŒÝ/°´ÀÄÈWÑUÔUˆØ�6˜7 FÐ*Ð*r=   c               ó  ‡!‡"‡#— |                      d| j        ¦  «        pt          j        j        |d<   t          j                             ¦   «         }t          j                             |¦  «        }| j        D ]
} ||i |¤Ž Œ| j	        |         \  }}	Š"}
 |
|i |¤Ž\  Š#}}t          |¦  «        t          |¦  «        z   }|                      |d ¦  «        }|�€y‰"                     |¦  «        }d„ | j        D ¦   «         }d„ |D ¦   «         }d„ t          ||¦  «        D ¦   «         }d|vs
J d¦   «         ‚d|vs
J d¦   «         ‚d	|vs
J d
¦   «         ‚|D ]!}||j        vr||vrt          d|z  ¦  «        ‚Œ"t!          |d„ ¦  «        }ˆ#fd„|D ¦   «         }d„ |D ¦   «         Š!t!          ‰!d„ ¦  «        }ˆ!ˆ"fd„|D ¦   «         }|                      t          j        j        ||||||g|¦  «        rd S |                      | |||¦  «        }|                      ||	|j        ¬¦  «        }|||<   |                      t          j        j        ||||||g|¦  «         t-          ¦   «         }| j                             ¦   «         D ]?\  \  }}\  }}|                      ||¦  «        x}|k    rt3          d|› d|› d|› �¦  «        ‚Œ@|sÆ|€J ‚t5          |¦  «        r |‰#¦  «        }t7          |¦  «        }|d         }|dk    r|d         nd}|dk    r|d         nd} |j        ||g‰#                     ¦   «         ¢R Ž }  |j        |||||j        |j         | t          j        j!        t          j        j"        g	‰#                     ¦   «         ¢R Ž  |S )NÚdebugc                ó   — g | ]	}|j         ‘Œ
S rp   rW  rð   s     r<   rt   z#JITFunction.run.<locals>.<listcomp><  s   € Ð3Ð3Ð3 !�q”vÐ3Ð3Ð3r=   c                ó   — g | ]
}|d          ‘ŒS rï   rp   rð   s     r<   rt   z#JITFunction.run.<locals>.<listcomp>=  s   € Ð4Ð4Ð4 �q˜”tÐ4Ð4Ð4r=   c                ó   — i | ]\  }}||“Œ	S rp   rp   )rr   rX   Úvs      r<   r  z#JITFunction.run.<locals>.<dictcomp>>  s   € ÐBÐBÐB¡& 1 a˜˜AÐBÐBÐBr=   Údevice_typez=device_type option is deprecated; current target will be usedrY  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Î   rp   )rZ   rb   s     r<   rí   z!JITFunction.run.<locals>.<lambda>G  s   € ¸sÀkÒ?Q€ r=   c           	     óp   •— i | ]2}|t          t          ‰                     ¦   «         ¦  «        |¦  «        “Œ3S rp   )r   r'   rÊ   )rr   ÚpathÚ
bound_argss     €r<   r  z#JITFunction.run.<locals>.<dictcomp>H  s;   ø€ ÐjÐjÐjÐW[˜$Õ 1µ$°z×7HÒ7HÑ7JÔ7JÑ2KÔ2KÈTÑ RÔ RÐjÐjÐjr=   c                ó   — g | ]
}|d          ‘ŒS rò   rp   rð   s     r<   rt   z#JITFunction.run.<locals>.<listcomp>J  s   € Ð5Ð5Ð5 ˜˜!œÐ5Ð5Ð5r=   c                ó,   — t          |t          ¦  «        S r?   )r,   rW   )rZ   ræ   s     r<   rí   z!JITFunction.run.<locals>.<lambda>K  s   € ½ÀAÅsÑ9KÔ9K€ r=   c           	     óX   •— i | ]&}|‰                      t          ‰|¦  «        ¦  «        “Œ'S rp   )Ú
parse_attrr   )rr   rX   ÚattrvalsrC  s     €€r<   r  z#JITFunction.run.<locals>.<dictcomp>L  s4   ø€ ÐZÐZÐZÈq�Q˜×*Ò*Õ+<¸XÀqÑ+IÔ+IÑJÔJÐZÐZÐZr=   )r–   r#  rO   z1 has changed since we compiled this kernel, from z to r   r™   r   )#ra   rz  r   Úruntimer   ru  Úget_current_deviceÚget_current_streamrn  Údevice_cachesrW   Úparse_optionsrg  r<  r&  ÚKeyErrorr   rk  Újit_cache_hookrr  rd  Újit_post_compile_hookÚobjectr7   r$  rS   rm  r$   Úlaunch_metadatarÊ   r  ÚfunctionÚpacked_metadataÚlaunch_enter_hookÚlaunch_exit_hook)$r9   r  r  rƒ   r  rY  r€  rh  Úkernel_cacher–   rw  rD  r#  rú   ÚkernelÚsigkeysÚsigvalsr  rX   Ú
constexprsr)  r:   Únot_presentr/   rZ   rb   Úglobals_dictÚnewValÚ	grid_sizeÚgrid_0Úgrid_1Úgrid_2r“  r‰  rC  r„  s$                                    @@@r<   r  zJITFunction.run#  sh  øøø€ Ø Ÿ*š* W¨d¬jÑ9Ô9ÐP½U¼]Ô=Pˆˆw‰õ ”×1Ò1Ñ3Ô3ˆÝ”×1Ò1°&Ñ9Ô9ˆð Ô&ð 	"ð 	"ˆDØˆD�$Ð!˜&Ð!Ð!Ð!Ð!à04Ô0BÀ6Ô0JÑ-ˆ�f˜g vð /5¨f°dÐ.E¸fÐ.EÐ.EÑ+ˆ
�N Gõ �.Ñ!Ô!¥C¨¡L¤LÑ0ˆØ×!Ò! # tÑ,Ô,ˆð ‰>à×+Ò+¨FÑ3Ô3ˆGà3Ð3 t¤{Ð3Ñ3Ô3ˆGØ4Ð4 ^Ð4Ñ4Ô4ˆGØBÐB­C°¸Ñ,AÔ,AÐBÑBÔBˆIà ¨Ð.Ð.Ð.Ð0oÑ.Ô.Ð.Ø 6Ð)Ð)Ð)Ð+eÑ)Ô)Ð)Ø 6Ð)Ð)Ð)Ð+eÑ)Ô)Ð)Øð ]ð ]�Ø˜GÔ,Ð,Ð,°¸'Ð1AÐ1AÝ"Ð#WÐZ[Ñ#[Ñ\Ô\Ð\øå& wÐ0QÐ0QÑRÔRˆJØjÐjÐjÐjÐ_iÐjÑjÔjˆJà5Ð5 nÐ5Ñ5Ô5ˆHÝ! (Ð,KÐ,KÑLÔLˆEØZÐZÐZÐZÐZÐTYÐZÑZÔZˆEØ�Š�uœ}Ô;¸SÀ)ÈVÐU_ÐahÐkpÐjqØ%ñ'ô 'ð à�tà—.’.  y°*¸eÑDÔDˆCØ—\’\ #¨f¸gÔ>N�\ÑOÔOˆFØ &ˆL˜ÑØ�OŠO�EœMÔ?ÀÀiÐQWÐYcÐelÐotÐnuØ"ñ$ô $ð $õ ‘h”hˆØ.2Ô.C×.IÒ.IÑ.KÔ.Kð 	qð 	qÑ*‰IˆT�1Ñ*˜˜\Ø&×*Ò*¨4°Ñ=Ô=Ð=�À#ÒEÐEÝ"Øo tÐoÐoÐ^aÐoÐoÐgmÐoÐoñqô qð qð Fð ð 	nàÐ#Ð#Ð#Ý˜‰~Œ~ð (Ø�t˜JÑ'Ô'�Ý˜D™	œ	ˆIØ˜!”WˆFØ )¨A¢ �T˜!”W�W°1ˆFØ )¨A¢ �T˜!”W�W°1ˆFà4˜fÔ4°T¸6ÐXÀJ×DUÒDUÑDWÔDWÐXÐXÐXˆOØˆFŒJ�v˜v v¨v°v´ÈÔH^Ð`oÝ”}Ô6½¼Ô8VðnØYc×YjÒYjÑYlÔYlðnð nð nð nàˆr=   c                óH   — | j         €| j        n|                       |¦  «        S r?   )Ú_reprÚ_fn_name)r9   rZ   s     r<   rc  zJITFunction.reprm  s   € Ø $¤
Ð 2ˆtŒ}ˆ}¸¿
º
À1¹¼ÐEr=   Nc	           	     óZ  — |r|ng }|r|ng }|| _         |j        | _        || _        t	          j        |¦  «        | _        || _        || _        t	          j        |¦  «        d         | _	        || _
        t          |¦  «        | _        || _        g | _        t          | j        j                             ¦   «         ¦  «        D ]I\  }	}
|	|v p|
j        |v }|	|v p|
j        |v }| j                             t)          |	|
||¦  «        ¦  «         ŒJt+          j        t	          j        |¦  «        ¦  «        }|t1          j        d|t0          j        ¦  «                             ¦   «         d …         }|                      |¦  «         t;          | j        ¦  «        | _        d | _         i | _!        d | _"        || _#        || _$        d„ | j        D ¦   «         | _%        d„ | j        D ¦   «         | _&        g | _'        |j(        | _(        |j)        | _)        |j*        | _*        |j+        | _+        |j        | _        d S )Nr™   z^def\s+\w+\s*\(c                ó   — g | ]	}|j         ‘Œ
S rp   rW  ©rr   Úps     r<   rt   z(JITFunction.__init__.<locals>.<listcomp>¡  s   € Ð6Ð6Ð6 Q˜!œ&Ð6Ð6Ð6r=   c                ó*   — g | ]}|j         ¯	|j        ‘ŒS rp   )rÐ   r¼   r©  s     r<   rt   z(JITFunction.__init__.<locals>.<listcomp>¢  s!   € ÐHÐHÐH Q¸¼ÐH˜1œ5ÐHÐHÐHr=   ),rL  rD   rL   ÚversionrF   r  r¿   rÁ   ÚgetsourcelinesÚstarting_line_numberr¥  rM  r¦  r“  rg  Ú	enumerater;  rÊ   r/   r	  r»   ÚtextwrapÚdedentÚ	getsourceÚreÚsearchÚ	MULTILINEÚstartÚ_unsafe_update_srcr   rx  r�  Úhashr7   r™  rz  rP   Ú	arg_namesrœ  rn  r¤   rT   r£   Ú__globals__)r9   rL  r¬  r¿   rÁ   rz  rP   rc  r“  Úir½   ÚdnsÚdns_oar:   s                 r<   r.   zJITFunction.__init__p  s)  € à1BÐJÐ-Ð-ÈÐØKiÐ)qÐ)GÐ)GÐoqÐ&àˆŒØ”mˆŒØˆŒÝ Ô*¨2Ñ.Ô.ˆŒØ!2ˆÔØ.LˆÔ+Ý$+Ô$:¸2Ñ$>Ô$>¸qÔ$AˆÔ!ØˆŒ
Ý% bÑ)Ô)ˆŒØ.ˆÔàˆŒÝ! $¤.Ô";×"BÒ"BÑ"DÔ"DÑEÔEð 	Cð 	C‰HˆAˆuØÐ(Ð(ÐK¨E¬JÐ:KÐ,KˆCØÐ8Ð8Ðh¸E¼JÐJhÐ<hˆFØŒK×Ò�{¨1¨e°S¸&ÑAÔAÑBÔBÐBÐBõ Œo�gÔ/°Ñ3Ô3Ñ4Ô4ˆØ•"”)Ð.°µR´\ÑBÔB×HÒHÑJÔJÐKÐKÔLˆØ×Ò Ñ$Ô$Ð$å(¨Ô);Ñ<Ô<ˆÔØˆŒ	ð TVˆÔð ˆŒØˆŒ
Ø ˆŒð 7Ð6¨$¬+Ð6Ñ6Ô6ˆŒØHÐH¨$¬+ÐHÑHÔHˆŒð  ˆÔð ”zˆŒØœˆŒØœOˆÔØœ>ˆÔØœ-ˆŒˆˆr=   c                óN   — | j         t          j        | j        ¦  «        j        z  S r?   )rº  rF   ÚgetclosurevarsrL  r5   rA   s    r<   Úget_capture_scopezJITFunction.get_capture_scope®  s    € ØÔ¥'Ô"8¸¼Ñ"AÔ"AÔ"KÑKÐKr=   c                óž  — | j         €Àt          j        | j        ¦  «        j        }t          | j        | j        || j        ¬¦  «        }| 	                    |  
                    ¦   «         ¦  «         |j        t          | j        ¦  «        z   | _         t          t          |j                             ¦   «         ¦  «        ¦  «        | _        | j         S )N)r/   r4   r5   r:   )r¸  rF   r¿  rL  r5   r   r¦  rº  r:   rq   ÚparserB   rW   r®  ÚdictÚsortedr7   r$  )r9   r5   Údependencies_finders      r<   rV   zJITFunction.cache_key±  s±   € ð Œ9ÐÝÔ.¨t¬wÑ7Ô7ÔAˆIÝ"4¸$¼-ÐQUÔQaÐmvØ9=¼ð#Cñ #Cô #CÐà×%Ò% d§j¢j¡l¤lÑ3Ô3Ð3Ø+Ô/µ#°dÔ6OÑ2PÔ2PÑPˆDŒIÝ$(­Ð0CÔ0T×0ZÒ0ZÑ0\Ô0\Ñ)]Ô)]Ñ$^Ô$^ˆDÔ!ØŒyÐr=   c                ó   — ddl m} |S )Nr   rÙ   )r®   rÏ   )r9   rÏ   s     r<   re   zJITFunction.type½  s   € à2Ð2Ð2Ð2Ð2Ð2ØÐr=   c               óR   —  | j         t          t          j        |¦  «        |ddœ|¤ŽS )NTr  )r  r>  Ú
MockTensorÚ
wrap_dtype)r9   r  rƒ   r  s       r<   r  zJITFunction.warmupÂ  s.   € ØˆtŒxµµZÔ5JÈDÑ1QÔ1Q˜T¨$Ð\Ð\ÐU[Ð\Ð\Ð\r=   c                óú  ‡— ddl m}m} dd l}dd lmŠ t          j                             ¦   «         }| 	                    |¦  «        }|d         | j
        k    r t          d|d         › d| j
        › �¦  «        ‚t          t          |d         ¦  «        }|d         }ˆfd	„t          ||¦  «        D ¦   «         }	t          t          |d
         ¦  «        }
|d         }t          t          |
|¦  «        ¦  «        }t          |d                              ¦   «         ¦  «        } || ||	|¦  «        }d„ |d                              ¦   «         D ¦   «         }|d         } ||d |¦  «        }|| j        |         d         |<   |S )Nr   )rd  rr  r   r/   zSpecialization data is for z but trying to preload for r  r   c                óz   •— i | ]7\  }}|‰j                              |¦  «        r‰                      |¦  «        n|“Œ8S rp   )r·   Úis_dtype)rr   rú   ry   Útls      €r<   r  z'JITFunction.preload.<locals>.<dictcomp>Ð  sR   ø€ ð 
ð 
ð 
á��Uð  B¤H×$5Ò$5°eÑ$<Ô$<ÐG�—’˜%‘”�À%ð
ð 
ð 
r=   r!  r"  r  c                ób   — i | ],\  }}|t          |t          ¦  «        rt          |¦  «        n|“Œ-S rp   )r,   r'   rë   r  s      r<   r  z'JITFunction.preload.<locals>.<dictcomp>Ù  sG   € ð 
ð 
ð 
á��Uð ¥¨Eµ4Ñ!8Ô!8ÐC•�u‘”�¸eð
ð 
ð 
r=   r#  rú   )rt  rd  rr  r%  Útriton.languager¯   r   ru  r‹  Úloadsr¦  rS   r>  rë   r<  rÃ  r$  r�  )r9   ra  rd  rr  r%  rY  Údeserialized_objr  r   r(  r!  r"  r)  r  r:   r#  rú   r™  rÍ  s                     @r<   ÚpreloadzJITFunction.preloadÅ  sÇ  ø€ Ø1Ð1Ð1Ð1Ð1Ð1Ð1Ð1ØˆˆˆØ$Ð$Ð$Ð$Ð$Ð$Ý”×1Ò1Ñ3Ô3ˆØŸ:š:Ð&9Ñ:Ô:ÐØ˜FÔ# t¤}Ò4Ð4ÝØrÐ.>¸vÔ.FÐrÐrÐcgÔcpÐrÐrñtô tð tå�EÐ#3°OÔ#DÑEÔEˆØ(¨Ô9ˆð
ð 
ð 
ð 
å! -°Ñ?Ô?ð
ñ 
ô 
ˆ	õ �Ð 0°Ô >Ñ?Ô?ˆ
Ø% lÔ3ˆ
Ý•S˜ ZÑ0Ô0Ñ1Ô1ˆÝÐ)¨+Ô6×<Ò<Ñ>Ô>Ñ?Ô?ˆ	Øˆi˜˜i¨°EÑ:Ô:ˆð
ð 
à.¨yÔ9×?Ò?ÑAÔAð
ñ 
ô 
ˆð ˜uÔ%ˆØ�˜˜d GÑ,Ô,ˆØ-3ˆÔ˜6Ô" 1Ô% cÑ*Øˆr=   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   )rg   rÂ  r:   r,   ÚModuler$   ÚbodyÚFunctionDef)r9   Útrees     r<   rÂ  zJITFunction.parseå  sg   € ÝŒy˜œÑ"Ô"ˆÝ˜$¥¤
Ñ+Ô+Ð+Ð+Ð+Ý�4”9‰~Œ~ Ò"Ð"Ð"Ð"Ý˜$œ) Aœ,­¬Ñ8Ô8Ð8Ð8Ð8Øˆr=   c                ó    — t          d¦  «        ‚)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rS   )r9   rƒ   r  s      r<   Ú__call__zJITFunction.__call__ì  s   € ÝÐWÑXÔXÐXr=   c                óŒ   •— |dk    rt          d|› d�¦  «        ‚t          t          | ¦  «                             ||¦  «         d S )Nr:   zCannot set attribute 'zX' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)ÚAttributeErrorr-   rQ   Ú__setattr__)r9   r/   ry   r;   s      €r<   rÜ  zJITFunction.__setattr__ï  s\   ø€ à�5Š=ˆ=Ý ð "-¸$ð "-ð "-ð "-ñ .ô .ð .õ 	�k˜4Ñ Ô ×,Ò,¨T°5Ñ9Ô9Ð9Ð9Ð9r=   c                óZ   •— d| _         t          ¦   «                              d|¦  «         dS )zŽ
        The only method allowed to modify src.
        Bypasses the __setattr__ restriction by calling super().__setattr__ directly.
        Nr:   )r¸  r-   rÜ  )r9   Únew_srcr;   s     €r<   r·  zJITFunction._unsafe_update_src÷  s,   ø€ ð
 ˆŒ	Ý‰Œ×Ò˜E 7Ñ+Ô+Ð+Ð+Ð+r=   c                ó2   — d| j         › d| j        j        › d�S )NzJITFunction(ú:r.  )rL   rL  r£   rA   s    r<   Ú__repr__zJITFunction.__repr__ÿ  s"   € ØC˜dœkÐCÐC¨D¬GÔ,@ÐCÐCÐCÐCr=   )r   rT  )NNNNNNN)rT   rD   r£   rS  rk  ro  rx  r  rc  r.   rÀ  r¥   rV   re   r  rÒ  rÂ  rÙ  rÜ  r·  rá  r¦   r§   s   @r<   rQ   rQ   Û  sx  ø€ € € € € ðð ð ð,
ð ,
ð ,
ð ,
ð\(ð (ð (ð+ð +ð +ðHð Hð HðTFð Fð Fð mqØ;?ð<(ð <(ð <(ð <(ð|Lð Lð Lð ð	ð 	ñ „Xð	ð ðð ñ „Xðð]ð ]ð ]ðð ð ð@ð ð ðYð Yð Yð:ð :ð :ð :ð :ð,ð ,ð ,ð ,ð ,ðDð Dð Dð Dð Dð Dð Dr=   rQ   rL  úJITFunction[T]c                ó   — d S r?   rp   rK  s    r<   Újiträ    s   € à€Cr=   ©r¬  rc  r“  r¿   rÁ   rz  rP   rc  úOptional[Callable]r“  r¿   úOptional[Iterable[int | str]]rÁ   rz  úOptional[bool]rP   úCallable[[T], JITFunction[T]]c                ó   — d S r?   rp   rå  s          r<   rä  rä    s	   € ð €Cr=   úOptional[T]ú4Union[JITFunction[T], Callable[[T], JITFunction[T]]]c               óF   ‡‡‡‡‡‡‡— 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
    rL  r   r   râ  c           
     ó¨   •— t          | ¦  «        sJ ‚t          j        j        rddlm}  || ‰‰‰‰‰‰‰¬¦  «        S t          | ‰‰‰‰‰‰‰¬¦  «        S )Nr™   )ÚInterpretedFunction)r¬  r¿   rÁ   rz  rP   rc  r“  )rm  r   rŠ  Ú	interpretÚinterpreterrï  rQ   )	rL  rï  rz  r¿   rÁ   r“  rP   rc  r¬  s	     €€€€€€€r<   Ú	decoratorzjit.<locals>.decorator8  s˜   ø€ Ý˜‰|Œ|ÐÐÐÝŒ=Ô"ð 	Ø8Ð8Ð8Ð8Ð8Ð8Ø&Ð& r°7ÐN_ØFdÐlqØ08¸tÐUdðfñ fô fð fõ ØØØ"3Ø/MØØ!ØØ /ð	ñ 	ô 	ð 	r=   N©rL  r   r   râ  rp   )	rL  r¬  rc  r“  r¿   rÁ   rz  rP   rò  s	    ``````` r<   rä  rä    sb   øøøøøøø€ ð:ð ð ð ð ð ð ð ð ð ð ð ð& 
€~Øˆy˜‰}Œ}Ðð Ðr=   c                  óZ   — e Zd ZdZed„ ¦   «         Zd„ Zed„ ¦   «         Zed„ ¦   «         ZdS )rÈ  zr
    Can be used in place of real tensors when calling:
        kernel.warmup(MockTensor(torch.float32), ...)
    c                óZ   — | j         j        dk    r| j        dk    rt          | ¦  «        S | S )Nr·   Útorch)r;   rT   rD   rÈ  r€   s    r<   rÉ  zMockTensor.wrap_dtype]  s/   € àŒ=Ô! WÒ,Ð,°´À7Ò1JÐ1JÝ˜c‘?”?Ð"Øˆ
r=   c                ó   — || _         d S r?   r  )r9   r·   s     r<   r.   zMockTensor.__init__c  s   € ØˆŒ
ˆ
ˆ
r=   c                 ó   — dS ©Nr   rp   rp   r=   r<   rá   zMockTensor.data_ptrf  ó   € àˆqr=   c                 ó   — dS rù  rp   rp   r=   r<   Ú	ptr_rangezMockTensor.ptr_rangej  rú  r=   N)	rT   rD   r£   r¤   ÚstaticmethodrÉ  r.   rá   rü  rp   r=   r<   rÈ  rÈ  W  s~   € € € € € ðð ð
 ðð ñ „\ðð
ð ð ð ðð ñ „\ðð ðð ñ „\ðð ð r=   rÈ  c                  óL   — 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dS )ÚTensorWrapperc                ót   — || _         || _        |j        | _        |j        | _        | j        j        | _        d S r?   )r·   rö   ÚdatarY  Úshape)r9   rö   r·   s      r<   r.   zTensorWrapper.__init__q  s1   € ØˆŒ
ØˆŒ	Ø”IˆŒ	Ø”kˆŒØ”Y”_ˆŒ
ˆ
ˆ
r=   c                ó4   — | j                              ¦   «         S r?   )rö   rá   rA   s    r<   rá   zTensorWrapper.data_ptrx  s   € ØŒy×!Ò!Ñ#Ô#Ð#r=   c                ó    —  | j         j        |Ž S r?   )rö   Ústride)r9   rƒ   s     r<   r  zTensorWrapper.stride{  s   € ØˆtŒyÔ Ð&Ð&r=   r   rW   c                ó(   — d| j         › d| j        › d�S )NzTensorWrapper[rX  r.  )r·   rö   rA   s    r<   Ú__str__zTensorWrapper.__str__~  s   € Ø: ¤
Ð:Ð:¨d¬iÐ:Ð:Ð:Ð:r=   c                ó4   — | j                              ¦   «         S r?   )rö   Úelement_sizerA   s    r<   r	  zTensorWrapper.element_size�  s   € ØŒy×%Ò%Ñ'Ô'Ð'r=   c                óZ   — t          | j                             ¦   «         | j        ¦  «        S r?   )rÿ  rö   Úcpur·   rA   s    r<   r  zTensorWrapper.cpu„  s   € Ý˜TœYŸ]š]™_œ_¨d¬jÑ9Ô9Ð9r=   c                óD   — | j                              |j         ¦  «         d S r?   )rö   Úcopy_)r9   Úothers     r<   r  zTensorWrapper.copy_‡  s   € ØŒ	�Š˜œ
Ñ#Ô#Ð#Ð#Ð#r=   c                óZ   — t          | j                             ¦   «         | j        ¦  «        S r?   )rÿ  rö   Úcloner·   rA   s    r<   r  zTensorWrapper.cloneŠ  s    € Ý˜TœYŸ_š_Ñ.Ô.°´
Ñ;Ô;Ð;r=   c                ó\   — t          | j                             |¦  «        | j        ¦  «        S r?   )rÿ  rö   Útor·   )r9   rY  s     r<   r  zTensorWrapper.to�  s"   € Ý˜TœYŸ\š\¨&Ñ1Ô1°4´:Ñ>Ô>Ð>r=   c                ó\   — t          | j                             |¦  «        | j        ¦  «        S r?   )rÿ  rö   Ú	new_emptyr·   )r9   Úsizess     r<   r  zTensorWrapper.new_empty�  s$   € Ý˜TœY×0Ò0°Ñ7Ô7¸¼ÑDÔDÐDr=   Nr×   )rT   rD   r£   r.   rá   r  r  r	  r  r  r  r  r  rp   r=   r<   rÿ  rÿ  o  s´   € € € € € ð%ð %ð %ð$ð $ð $ð'ð 'ð 'ð;ð ;ð ;ð ;ð(ð (ð (ð:ð :ð :ð$ð $ð $ð<ð <ð <ð?ð ?ð ?ðEð Eð Eð Eð 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 rJ  )r,   rÿ  rö   r·   rê   r›   re   )râ   r·   s     r<   Úreinterpretr  ”  s„   € Ý�&�-Ñ(Ô(ð AØ�F”KÔ%Ò%Ð%à”;Ðõ ! ¤¨eÑ4Ô4Ð4Ý	�˜Ñ	$Ô	$ð Aå˜V UÑ+Ô+Ð+åÐ?µ°V±´Ð?Ð?Ð?Ñ@Ô@Ð@r=   c                óP  — | }t          |t          ¦  «        s|j        }t          |t          ¦  «        ¯|j        j        j        }t          j        |j        ¦  «        \  }}t          |¦  «        D ]3\  }}|                     ¦   «          	                    d¦  «        r||z  } nŒ4||fS )Nzdef )
r,   rQ   rL  Ú__code__Úco_filenamerF   r­  r¯  r±   rI   )rL  Úbase_fnÚ	file_nameÚlinesÚ
begin_lineÚidxÚlines          r<   Úget_jit_fn_file_liner!  £  s´   € Ø€GÝ˜¥+Ñ.Ô.ð Ø”*ˆõ ˜¥+Ñ.Ô.ð à”
Ô#Ô/€IÝÔ.¨w¬zÑ:Ô:Ñ€Eˆ:õ ˜uÑ%Ô%ð ð ‰	ˆˆTØ�:Š:‰<Œ<×"Ò" 6Ñ*Ô*ð 	Ø˜#ÑˆJØˆEð	ð �jÐ Ð r=   r×   )Fró  )rc  ræ  r“  ræ  r¿   rç  rÁ   rç  rz  rè  rP   rè  r   ré  r?   )rL  rë  rc  ræ  r“  ræ  r¿   rç  rÁ   rç  rz  rè  rP   rè  r   rì  )?Ú
__future__r   r   rg   rk   r0   rF   r‹   r³  r°  Úcollectionsr   Údataclassesr   Ú	functoolsr   Útypingr   r	   r
   r   r   r   r   r   r   r   Útriton.tools.tensor_descriptorr   Útypesr   rE   r   Úruntime.driverr   Ú_utilsr   r   r   r   rT   r$   rJ   r   ÚNodeVisitorr   r³   r»   rõ   r  r  r  r  r,  rH  rM  rO  rQ   rä  rÈ  rÿ  r  r!  rp   r=   r<   ú<module>r,     s  ðØ ,Ð ,Ð ,Ð ,Ð ,Ð ,Ð ,Ð ,Ø 
€
€
€
Ø €€€Ø €€€Ø €€€Ø Ð Ð Ð Ø 	€	€	€	Ø €€€Ø #Ð #Ð #Ð #Ð #Ð #Ø !Ð !Ð !Ð !Ð !Ð !Ø %Ð %Ð %Ð %Ð %Ð %Ø dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dÐ dà ;Ð ;Ð ;Ð ;Ð ;Ð ;Ø Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð Ø #Ð #Ð #Ð #Ð #Ð #Ø eÐ eÐ eÐ eÐ eÐ eÐ eÐ eÐ eÐ eÐ eÐ eàÐ.˜3˜3˜~Ñ.Ô.Ð.Ð.Ô/€à€GˆC�L„L€ðH!ð H!ð H!ð H!ð H!˜œñ H!ô H!ð H!ð`Dð Dð Dð Dð4/>ð />ð />ð />ð />ñ />ô />ð />ðd €	ØÐ ð6ð 6ð 6ðr@ð @ð @ð @ð	Zð 	Zð 	Zð 	Zð 	Z�g˜a”jñ 	Zô 	Zð 	Zð	ð 	ð 	ð7*ð 7*ð 7*ðt0ð 0ð 0ð ðð ð ð ð ñ ô ñ „ððeDð eDð eDð eDð eD�/ !Ô$ñ eDô eDð eDðZ	 
ðð ð ñ 
„ðð 
ð Ø#Ø*.Ø7;ØDHØ Ø#ð
ð 
ð 
ð 
ð 
ñ 
„ð
ð ð4ð Ø#Ø*.Ø7;ØDHØ Ø#ð4ð 4ð 4ð 4ð 4ð 4ðxð ð ð ð ñ ô ð ð0"Eð "Eð "Eð "Eð "Eñ "Eô "Eð "EðJAð Að Að!ð !ð !ð !ð !r=   