§
    ÕóÉ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
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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% d	dl&m'Z'm(Z(m)Z)m*Z* ddl+m,Z, d dl-m.Z.m/Z/m0Z0 dZ1dZ2 e	j3        de	j4        ¦  «        Z5 ed¦  «        Z6 ed¦  «        Z7 ed¦  «        Z8 ed¦  «        Z9 G d„ dej:        ¦  «        Z;dMd„Z< G d„ d¦  «        Z=dNd „Z> G d!„ d"ee6         ¦  «        Z?d#„ Z@d$„ ZAd%„ ZB G d&„ d'¦  «        ZCi ZDe G d(„ d)¦  «        ¦   «         ZEd*„ ZFd+„ ZG G d,„ d-eCe?e6         ¦  «        ZHedOd0„¦   «         ZIedddddddd1œdPd<„¦   «         ZI	 dQdddddddd1œdRd?„ZI G d@„ dA¦  «        ZJ G dB„ dC¦  «        ZKdD„ ZLdE„ ZMdF„ ZN G dG„ dHeC¦  «        ZO G dI„ dJeCee6         ¦  «        ZPdSdL„ZQdS )Té    )ÚannotationsÚdivisionN)Údefaultdict)Ú	dataclass)Úcached_property)ÚCallableÚConcatenateÚGenericÚIterableÚOptionalÚ	ParamSpecÚTYPE_CHECKINGÚTypeVarÚoverloadÚDictÚAnyÚTuple)ÚBaseBackend)Ú
ModuleTypeé   )Úknobsé   )Údriver)Ú_async_compile)Úfind_paths_ifÚget_iterable_pathÚtype_canonicalisation_dictÚis_namedtuple)Úget_cache_key)Úget_cache_invalidating_env_varsÚnative_specialize_implÚirztriton.languagez"triton.experimental.gluon.languagez!^(?P<indent>[ \t]*)def\s+\w+\s*\(ÚTÚPÚRÚUc                  ó~   ‡ — e Zd ZdZdˆ fd„Zed„ ¦   «         Zd„ Zd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£| _	        t          t          ddh| _        i | _        d| _        d S )Núutf-8>
   ÚintÚlenÚmaxÚminÚlistÚfloatÚprintÚrangeÚgetattrÚ
isinstanceÚcopyÚmathF)ÚsuperÚ__init__ÚnameÚhashlibÚsha256ÚencodeÚhasherÚglobalsÚ	nonlocalsÚsupported_python_builtinsÚGLUON_MODULEÚTRITON_MODULEÚsupported_modulesÚused_global_valsÚvisiting_arg_default_value)Úselfr;   r@   rA   ÚsrcÚ	__class__s        €úw/home/lighthouse/badge-speaker-benchmark/.benchmark-venv-community/lib64/python3.11/site-packages/triton/runtime/jit.pyr:   zDependenciesFinder.__init__3   s“   ø€ Ý‰Œ×ÒÑÔÐØˆŒ	Ý”n S§Z¢Z°Ñ%8Ô%8Ñ9Ô9ˆŒð ˆŒØ"ˆŒð*
ð *
ð *
ˆÔ&õ ÝØØð	"
ˆÔð" TVˆÔà*/ˆÔ'Ð'Ð'ó    c                ó4   — | j                              ¦   «         S ©N)r?   Ú	hexdigest©rH   s    rK   ÚretzDependenciesFinder.ret^   s   € àŒ{×$Ò$Ñ&Ô&Ð&rL   c                ó0  — t          |t          ¦  «        sJ ‚|j        }| j                             ¦   «         |j                             ¦   «         z  D ]V}|\  }}| j        |         \  }}|j        |         \  }}||k    r)t          d|› d|› d| j        › d|j        › d|› d�¦  «        ‚ŒW| j                             |j        ¦  «         |t          t          |dd¦  «        ¦  «        z  }| j                             |                     d	¦  «        ¦  «         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,   )r6   ÚJITCallableÚ	cache_keyrF   ÚkeysÚRuntimeErrorr;   Ú__name__ÚupdateÚstrr5   r?   r>   )rH   ÚfuncÚfunc_keyÚkÚvar_nameÚ_Úv1Úv2s           rK   Ú_update_hashzDependenciesFinder._update_hashb   s_  € Ý˜$¥Ñ,Ô,Ð,Ð,Ð,Ø”>ˆð Ô&×+Ò+Ñ-Ô-°Ô0E×0JÒ0JÑ0LÔ0LÑLð 	ð 	ˆAØ‰KˆH�aØÔ)¨!Ô,‰EˆB�ØÔ)¨!Ô,‰EˆB�Ø�RŠxˆxÝ"ð P xð  Pð  P¸Bð  Pð  PÐPTÔPYð  Pð  PÐnrÔn{ð  Pð  Pð  UWð  Pð  Pð  Pñô ð ð ð 	Ô×$Ò$ TÔ%:Ñ;Ô;Ð;à•C�  j°%Ñ8Ô8Ñ9Ô9Ñ9ˆØŒ×Ò˜8Ÿ?š?¨7Ñ3Ô3Ñ4Ô4Ð4Ð4Ð4rL   Nc                óÐ  — ddl m} |�t          |¦  «        t          u rd S t	          |dd¦  «        r`| j                             t          |j        ¦  «         	                    d¦  «        ¦  «         |j
        D ]}|                      |¦  «         Œd S t	          |dd¦  «        rd S t	          |dd¦  «        d	k    rd S t          |t          ¦  «        r|                      |¦  «         d S t          |¦  «        r7t          |t          ¦  «        s"t          ||¦  «        st!          d
|› �¦  «        ‚| j        rd S |�-t%          j        |¦  «        |f| j        |t+          |¦  «        f<   d S )Nr   ©Ú	constexprÚ__triton_aggregate__Fr,   Ú__triton_builtin__Ú
__module__Ú ztriton.language.extra.libdevicez!Unsupported function referenced: )Úlanguage.corerf   Útyper   r5   r?   rZ   r[   Ú__annotations__r>   Ú
hash_attrsÚrecord_referencer6   rU   rc   ÚcallablerX   rG   r7   ÚdeepcopyrF   Úid)rH   ÚvalÚvar_dictr;   rf   Úattrs         rK   ro   z#DependenciesFinder.record_referencet   s“  € Ø-Ð-Ð-Ð-Ð-Ð-ð ˆ;�$˜s™)œ)¥zÐ1Ð1ØˆFå�3Ð.°Ñ6Ô6ð 	ØŒK×Ò�s 3Ô#6Ñ7Ô7×>Ò>¸wÑGÔGÑHÔHÐHØœð ,ð ,�Ø×%Ò% dÑ+Ô+Ð+Ð+ØˆFå�3Ð,¨eÑ4Ô4ð 	ØˆFõ �3˜ bÑ)Ô)Ð-NÒNÐNØˆFå�c�;Ñ'Ô'ð 	Ø×Ò˜cÑ"Ô"Ð"ØˆFå�C‰=Œ=ð 	J¥¨CµÑ!6Ô!6ð 	J½zÈ#ÈyÑ?YÔ?Yð 	JÝÐHÀ3ÐHÐHÑIÔIÐIð Ô*ð 	ØˆFàÐÝ;?¼=ÈÑ;MÔ;MÈxÐ:XˆDÔ! 4­¨H©¬Ð"6Ñ7ØˆrL   c                óþ   ‡ — t          |j        ¦  «        t          j        u r|j        S |j        ‰ j        v rd S ˆ fd„} ||j        ¦  «        \  }}|j        ‰ j        v r|S ‰                      |||j        ¦  «         |S )Nc                ó    •— ‰j                              | d ¦  «        }|�	|‰j         fS ‰j                             | d ¦  «        }|�	|‰j        fS dS )N©NN)r@   ÚgetrA   )r;   rs   rH   s     €rK   Úname_lookupz2DependenciesFinder.visit_Name.<locals>.name_lookup¤   sZ   ø€ Ø”,×"Ò" 4¨Ñ.Ô.ˆCØˆØ˜DœLÐ(Ð(Ø”.×$Ò$ T¨4Ñ0Ô0ˆCØˆØ˜DœNÐ*Ð*Ø�:rL   )rl   ÚctxÚastÚStorerr   Úlocal_namesrB   ro   )rH   Únoderz   rs   rt   s   `    rK   Ú
visit_NamezDependenciesFinder.visit_Nameœ   s˜   ø€ Ý�”‰>Œ>�SœYÐ&Ð&Ø”7ˆNàŒ7�dÔ&Ð&Ð&à�4ð	ð 	ð 	ð 	ð 	ð $˜ D¤GÑ,Ô,‰ˆˆXØŒ7�dÔ4Ð4Ð4ØˆJà×Ò˜c 8¨T¬WÑ5Ô5Ð5Øˆ
rL   c                ó*   ‡ — ˆ fd„|j         D ¦   «         S )Nc                ó:   •— g | ]}‰                      |¦  «        ‘ŒS © ©Úvisit)Ú.0ÚeltrH   s     €rK   ú
<listcomp>z2DependenciesFinder.visit_Tuple.<locals>.<listcomp>·   s#   ø€ Ð5Ð5Ð5 C�—
’
˜3‘”Ð5Ð5Ð5rL   )Úelts©rH   r   s   ` rK   Úvisit_TuplezDependenciesFinder.visit_Tuple´   s!   ø€ ð 6Ð5Ð5Ð5¨4¬9Ð5Ñ5Ô5Ð5rL   c                óh  — |                       |j        ¦  «        }t          |t          j        ¦  «        r4|                       |j        ¦  «        }t          |t          j        ¦  «        °4t          |dd¦  «        }|�	|| j        v rd S t          ||j        d ¦  «        }|                      |¦  «         |S )NrY   rj   )	r…   Úvaluer6   r|   Ú	Attributer5   rE   ru   ro   )rH   r   ÚlhsÚlhs_namerQ   s        rK   Úvisit_Attributez"DependenciesFinder.visit_Attribute¹   s¦   € Ø�jŠj˜œÑ$Ô$ˆÝ˜�cœmÑ,Ô,ð 	(Ø—*’*˜SœYÑ'Ô'ˆCõ ˜�cœmÑ,Ô,ð 	(å˜3 
¨BÑ/Ô/ˆØˆ;˜( dÔ&<Ð<Ð<Ø�4Ý�c˜4œ9 dÑ+Ô+ˆØ×Ò˜cÑ"Ô"Ð"Øˆ
rL   c                óf   — d„ |j         j         D ¦   «         | _        |                      |¦  «         d S )Nc                ó   — h | ]	}|j         ’Œ
S rƒ   ©Úarg)r†   r•   s     rK   ú	<setcomp>z7DependenciesFinder.visit_FunctionDef.<locals>.<setcomp>Æ   s   € Ð>Ð>Ð>¨˜CœGÐ>Ð>Ð>rL   )Úargsr~   Úgeneric_visitrŠ   s     rK   Úvisit_FunctionDefz$DependenciesFinder.visit_FunctionDefÄ   s6   € à>Ð>¨t¬y¬~Ð>Ñ>Ô>ˆÔØ×Ò˜4Ñ Ô Ð Ð Ð rL   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)rG   r…   )ÚdefaultsÚexprrH   s     €rK   Úvisit_defaultsz:DependenciesFinder.visit_arguments.<locals>.visit_defaultsÕ   sp   ø€ ð8ØÔ:Ð:Ð:Ð:Ø26�Ô/Ø$ð )ð )�DØÐ'ØŸ
š
 4Ñ(Ô(Ð(øð)ð 38�Ô/Ð/Ð/ø°%�Ô/Ð7Ð7Ð7Ð7s	   ƒ,9 ¹	A)
Ú	itertoolsÚchainÚposonlyargsr—   ÚvarargÚ
kwonlyargsr…   Úkw_defaultsÚkwargrœ   )rH   r   rž   r•   s   `   rK   Ú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”}Ñ%Ô%Ð%Ð%Ð%rL   c                óV   ‡ ‡— ‰                       |¦  «        }ˆˆ fd„Š ‰|¦  «         d S )Nc                óŠ   •— t          | t          ¦  «        r| D ]} ‰|¦  «         Œd S ‰j                             | ¦  «         d S rN   )r6   r1   r~   Úadd)ÚtÚsubÚ_addrH   s     €€rK   r¬   z0DependenciesFinder.visitAssnTarget.<locals>._addí   s\   ø€ Ý˜!�TÑ"Ô"ð (Øð ð �CØ�D˜‘I”I�I�Iðð ð Ô ×$Ò$ QÑ'Ô'Ð'Ð'Ð'rL   r„   )rH   r   Útargetr¬   s   `  @rK   ÚvisitAssnTargetz"DependenciesFinder.visitAssnTargeté   sF   øø€ à—’˜DÑ!Ô!ˆð	(ð 	(ð 	(ð 	(ð 	(ð 	(ð 	ˆˆV‰ŒˆˆˆrL   c                ó¾   — t          |j        ¦  «        dk    rt          d¦  «        ‚|                      |j        d         ¦  «         |                      |¦  «         d S )Nr   z2Simultaneous multiple assignment is not supported.r   )r.   ÚtargetsÚ	TypeErrorr®   r˜   rŠ   s     rK   Úvisit_AssignzDependenciesFinder.visit_Assignö   s^   € ÝˆtŒ|ÑÔ Ò!Ð!õ
 ÐPÑQÔQÐQà×Ò˜Tœ\¨!œ_Ñ-Ô-Ð-ð 	×Ò˜4Ñ Ô Ð Ð Ð rL   c                ód   — |                       |j        ¦  «         |                      |¦  «         d S rN   ©r®   r­   r˜   rŠ   s     rK   Úvisit_AnnAssignz"DependenciesFinder.visit_AnnAssign  ó4   € Ø×Ò˜Tœ[Ñ)Ô)Ð)ð 	×Ò˜4Ñ Ô Ð Ð Ð rL   c                ód   — |                       |j        ¦  «         |                      |¦  «         d S rN   r´   rŠ   s     rK   Ú	visit_ForzDependenciesFinder.visit_For	  r¶   rL   )r)   r*   rx   )rY   ri   Ú__qualname__Ú__doc__r:   ÚpropertyrQ   rc   ro   r€   r‹   r‘   r™   r¦   r®   r²   rµ   r¸   Ú__classcell__©rJ   s   @rK   r(   r(   '   s  ø€ € € € € ð	ð 	ð)0ð )0ð )0ð )0ð )0ð )0ðV ð'ð 'ñ „Xð'ð5ð 5ð 5ð$&ð &ð &ð &ðPð ð ð06ð 6ð 6ð
	ð 	ð 	ð!ð !ð !ð
&ð &ð &ð@ð ð ð!ð !ð !ð!ð !ð !ð!ð !ð !ð !ð !ð !ð !rL   r(   r)   r[   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.Ú_trj   )Útriton.language.coreÚlanguageÚcorer6   r[   ÚstripÚ
startswithÚremoveprefixÚ_normalize_tyÚendswithÚpointer_typeÚ
element_tyÚdtyper;   rl   rY   r   ry   Úreplace)ÚtyrÆ   s     rK   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ÐCrL   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 rN   )rÓ   Ú_paramrÖ   rØ   )rH   rÓ   rÔ   rÖ   rØ   s        rK   r:   zKernelParam.__init__2  s&   € àˆŒØˆŒØ!2ˆÔØ.LˆÔ+Ð+Ð+rL   c                ó   — | j         j        S rN   )rÚ   r;   rP   s    rK   r;   zKernelParam.name9  s   € àŒ{ÔÐrL   r)   r[   c                óŽ   — | j         j        r| j         j        t          j        j        k    rdS t          | j         j        ¦  «        S )Nrj   )rÚ   Ú
annotationÚinspectÚ	ParameterÚemptyrÊ   rP   s    rK   rÝ   zKernelParam.annotation=  s<   € àŒ{Ô%ð 	¨¬Ô)?Å7ÔCTÔCZÒ)ZÐ)ZØ�2Ý˜Tœ[Ô3Ñ4Ô4Ð4rL   c                óä   — | j         }|                     d¦  «        r|dd …         }n|                     d¦  «        r
|dd …         }|t          t          j        ¦   «         ¦  «        v r| j         S dS )NrÁ   r   rÀ   r   rj   )rÝ   rÈ   Úsetr   Úvalues)rH   Úas     rK   Úannotation_typezKernelParam.annotation_typeC  sw   € àŒOˆØ�<Š<˜ÑÔð 	Ø�!�"�"”ˆAˆAØ�\Š\˜#ÑÔð 	Ø�!�"�"”ˆAØ•Õ.Ô5Ñ7Ô7Ñ8Ô8Ð8Ð8Ø”?Ð"ØˆrrL   c                ó   — d| j         v S ©Nrf   )rÝ   rP   s    rK   Úis_constexprzKernelParam.is_constexprN  s   € à˜dœoÐ-Ð-rL   c                óZ   — | j         rdS d| j        v p| j                             d¦  «        S )NFr¿   rÁ   )rè   rÝ   rÈ   rP   s    rK   Úis_constzKernelParam.is_constR  s4   € àÔð 	Ø�5Ø˜$œ/Ð)ÐM¨T¬_×-GÒ-GÈÑ-MÔ-MÐMrL   c                ó   — | j         j        S rN   )rÚ   ÚdefaultrP   s    rK   rì   zKernelParam.defaultX  s   € àŒ{Ô"Ð"rL   c                ó@   — | j         j        t          j        j        k    S rN   )rÚ   rì   rÞ   rß   rà   rP   s    rK   Úhas_defaultzKernelParam.has_default\  s   € àŒ{Ô"¥gÔ&7Ô&=Ò=Ð=rL   N)rÓ   r-   rÔ   rÕ   rÖ   r×   rØ   r×   ©r)   r[   )rY   ri   r¹   rº   r:   r   r;   rÝ   rå   rè   rê   r»   rì   rî   rƒ   rL   rK   rÒ   rÒ   /  sü   € € € € € ØLÐLðMð Mð Mð Mð ð ð  ñ „_ð ð ð5ð 5ð 5ñ „_ð5ð
 ðð ð ñ „_ðð ð.ð .ñ „_ð.ð ðNð Nñ „_ðNð
 ð#ð #ñ „Xð#ð ð>ð >ñ „Xð>ð >ð >rL   rÒ   Fc                óF   — d}d}t          t          | |||¦  «        d         S )NFTr   )r!   r   )r•   Ú
specializerê   Úaligns       rK   Úmangle_typeró   a  s(   € Ø€HØ€EÝ!¥+¨s°H¸jÈ%ÑPÔPÐQRÔSÐSrL   c                  ó.   — e Zd ZU ded<   d„ Zd„ Zdd„ZdS )	ÚKernelInterfacer#   Úrunc               óR   —  | j         t          t          j        |¦  «        |ddœ|¤ŽS )NT©ÚgridÚwarmup)rö   ÚmapÚ
MockTensorÚ
wrap_dtype)rH   rù   r—   Úkwargss       rK   rú   zKernelInterface.warmupj  s.   € ØˆtŒxµµZÔ5JÈDÑ1QÔ1Q˜T¨$Ð\Ð\ÐU[Ð\Ð\Ð\rL   c               ó    — t          d¦  «        ‚)Nzrun not implemented)ÚNotImplementedError)rH   rù   rú   r—   rþ   s        rK   rö   zKernelInterface.runm  s   € Ý!Ð"7Ñ8Ô8Ð8rL   r)   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 )NFrø   )rö   )r—   rþ   rù   rH   s     €€rK   ú<lambda>z-KernelInterface.__getitem__.<locals>.<lambda>v  s    ø€  x t¤xÈ$°TÀ%Ð'YÐ'YÐRXÐ'YÐ'Y€ rL   rƒ   )rH   rù   s   ``rK   Ú__getitem__zKernelInterface.__getitem__p  s   øø€ ð ZÐYÐYÐYÐYÐYrL   N)r)   r#   )rY   ri   r¹   rm   rú   rö   r  rƒ   rL   rK   rõ   rõ   g  sZ   € € € € € € Ø
€F€F�Fð]ð ]ð ]ð9ð 9ð 9ðZð Zð Zð Zð Zð ZrL   rõ   c           
     óŒ  — d„ |                      ¦   «         D ¦   «         }dd l}| |d„ |                     ¦   «         D ¦   «         t          |                     ¦   «         ¦  «        d„ |                     ¦   «         D ¦   «         t          |                     ¦   «         ¦  «        |j        ||j        dœ	}|                     |¦  «        }	|	S )Nc                óÚ   — i | ]h\  }}||j         j        d k    rt          |¦  «        nB|j         j        dk    r	d|j        in)|j         j        dk    rd|j        › d|j        j        › �in|“ŒiS )rÎ   rf   ÚJITFunctionÚjit_functionú:)rJ   rY   r[   r�   ÚmoduleÚfnr¹   ©r†   Úkeyr�   s      rK   ú
<dictcomp>z1serialize_specialization_data.<locals>.<dictcomp>{  s¢   € ð ð ð ñ ˆC�ð 	˜5œ?Ô3°wÒ>Ð>�S�‰ZŒZˆZØŒ?Ô# {Ò2Ð2ð FQÐRWÔR]ÐD^ÐD^àŒ?Ô# }Ò4Ð4ð :HÈEÌLÐIrÐIrÐ[`Ô[cÔ[pÐIrÐIrÐ8sÐ8sØ:?ðð ð rL   r   c                ó,   — g | ]}t          |¦  «        ‘ŒS rƒ   ©r1   ©r†   Úxs     rK   rˆ   z1serialize_specialization_data.<locals>.<listcomp>„  s   € Ð?bÐ?bÐ?bÈAÅÀQÁÄÐ?bÐ?bÐ?brL   c                ó,   — g | ]}t          |¦  «        ‘ŒS rƒ   r  r  s     rK   rˆ   z1serialize_specialization_data.<locals>.<listcomp>…  s   € Ð0OÐ0OÐ0O¸Qµ°a±´Ð0OÐ0OÐ0OrL   )	r;   Ú	signatureÚconstant_keysÚconstant_valsÚ
attrs_keysÚ
attrs_valsÚoptionsr  r­   )ÚitemsÚjsonrW   r1   rã   Ú__dict__Údumps)
r;   r  Ú	constantsÚattrsr  r  r­   r  ÚobjÚserialized_objs
             rK   Úserialize_specialization_datar"  z  sÑ   € ðð ð $Ÿ/š/Ñ+Ô+ð	ñ ô €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¸6¼?ðð €Cð
 —Z’Z ‘_”_€NØÐrL   c           
     ój  — 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<   ||d<   t"          |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(backend, ú, Úu1Nr   )ÚfpÚbfFz("z",) + z[1:]z", None)c                óŒ   — | \  }}|j         t          j        j        k    rd|› �S |j        t          j        j        u r|n|› d|› �S )NrÀ   z	=default_)ÚkindrÞ   rß   ÚVAR_POSITIONALrì   rà   )Ú
name_paramr;   rÔ   s      rK   r•   z+create_function_from_signature.<locals>.arg­  sU   € Ø ‰ˆˆeØŒ:�Ô*Ô9Ò9Ð9Ø�t�:�:ÐØ”}­Ô(9Ô(?Ð?Ð?ˆtˆtÈÐE]ÐE]ÐW[ÐE]ÐE]Ð]rL   z
def dynamic_func(z	**optionsz):
    params = {c                ó   — g | ]
}d |› d|› �‘ŒS )ú'z': rƒ   )r†   r;   s     rK   rˆ   z2create_function_from_signature.<locals>.<listcomp>µ  s)   € ÐQÐQÐQ°4Ð.˜tÐ.Ð.¨Ð.Ð.ÐQÐQÐQrL   z}
    specialization = [ú,z-]
    return params, specialization, options
c                óZ   — i | ](\  }}|j         t          j        j        u¯d |› �|j         “Œ)S )Údefault_)rì   rÞ   rß   rà   )r†   r;   rÔ   s      rK   r  z2create_function_from_signature.<locals>.<dictcomp>»  sE   € ð ð ð áˆD�%ØŒ=¥Ô 1Ô 7Ð7Ð7ð 	�4ÐÐ˜5œ=à7Ð7Ð7rL   Úspecialize_implÚbackendrU   Údynamic_func)r.   Ú
parametersÚziprW   rè   Úappendrê   rÖ   rØ   rå   r6   r[   Újoinr1   rû   r  r!   rU   Úexec)ÚsigÚkparamsr5  Úspecializationr;   Úkprê   rñ   rò   rQ   r•   Ú	func_bodyÚfunc_namespacer4  s                 rK   Úcreate_function_from_signaturerB  Œ  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ØX¨dÐXÐX°hÐXÐXÀ*ÐXÐXÐPUÐXÐXÐXˆ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Ñ/Ô/Ð/Ð/ð^ð ^ð ^ðØ—)’)�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õ -€OØ(7€NÐ$Ñ%Ø '€N�9ÑÝ$/€N�=Ñ!õ 	ˆ�NÑ#Ô#Ð#ð ˜.Ô)Ð)rL   c                ó$   — | j         › d| j        › �S )Nú.)ri   r¹   ©r  s    rK   Úget_full_namerF  Í  s   € ØŒmÐ/Ð/˜bœoÐ/Ð/Ð/rL   c                  óˆ   — e Zd Zd„ Zd„ Zedd„¦   «         Zd„ Zd„ Zed„ ¦   «         Z	dd„Z
d„ Zd„ Zd„ Z eee¬¦  «        ZdS )rU   c                ó"  — || _         t          j        |¦  «        | _        	 t          j        |¦  «        \  | _        | _        n"# t          $ r}t          d¦  «        |‚d }~ww xY wt          |¦  «        | _	        t          j        ¦   «         | _        d                     | j        ¦  «        }|j        j        | _        t#          | j        | j        ¦  «        | _        t'          |¦  «        | _        t+          j        |¦  «        }|t/          j        d|t.          j        ¦  «                             ¦   «         d …         }|| _        d | _        i | _        |j        | _        |j        | _        |j         | _         |j!        | _!        |j"        | _"        d S )Nz1@jit functions should be defined in a Python filerj   z^def\s+\w+\s*\()#r  rÞ   r  ÚgetsourcelinesÚraw_srcÚstarting_line_numberÚOSErrorÚ
ValueErrorrF  Ú_fn_nameÚ	threadingÚRLockÚ
_hash_lockr:  Ú__code__Úco_filenameÚ	file_nameÚget_def_line_numberÚdef_file_line_numberÚget_def_col_numberÚdef_file_col_numberÚtextwrapÚdedentÚreÚsearchÚ	MULTILINEÚstartÚ_srcÚhashrF   rº   rY   r¹   Ú__globals__ri   )rH   r  ÚeÚraw_src_strrI   s        rK   r:   zJITCallable.__init__Ó  sf  € ØˆŒÝ Ô*¨2Ñ.Ô.ˆŒð	YÝ6=Ô6LÈRÑ6PÔ6PÑ3ˆDŒL˜$Ô3Ð3øÝð 	Yð 	Yð 	YÝÐPÑQÔQÐWXÐXøøøøð	Yøøøå% bÑ)Ô)ˆŒÝ#œ/Ñ+Ô+ˆŒð —g’g˜dœlÑ+Ô+ˆð œÔ0ˆŒÝ$7¸¼ÀdÔF_Ñ$`Ô$`ˆÔ!Ý#5°kÑ#BÔ#BˆÔ åŒo˜kÑ*Ô*ˆØ•"”)Ð.°µR´\ÑBÔB×HÒHÑJÔJÐKÐKÔLˆØˆŒ	ØˆŒ	ð TVˆÔð ”zˆŒØœˆŒØœOˆÔØœ>ˆÔØœ-ˆŒˆˆs   ¢!A Á
A#ÁAÁA#c                ó’   — | j         }|j        €| j        S d„ t          |j        j        |j        ¦  «        D ¦   «         }| j        |z  S )Nc                ó$   — i | ]\  }}||j         “ŒS rƒ   )Úcell_contents)r†   r;   Úcells      rK   r  z1JITCallable.get_capture_scope.<locals>.<dictcomp>   s!   € ÐmÐmÐm±*°$¸�T˜4Ô-ÐmÐmÐmrL   )r  Ú__closure__ra  r8  rR  Úco_freevars)rH   r  rA   s      rK   Úget_capture_scopezJITCallable.get_capture_scopeü  sN   € ØŒWˆØŒ>Ð!ØÔ#Ð#ØmÐmÅÀBÄKÔD[Ð]_Ô]kÑ@lÔ@lÐmÑmÔmˆ	ØÔ )Ñ+Ð+rL   r)   r[   c                óf  ‡— | j         5  | j        �| j        cd d d ¦  «         S d| j        › �| _        t          j        | j        ¦  «        j        }t          | j        | j        || j	        ¬¦  «        }| 
                    |                      ¦   «         ¦  «         |j        t          | j        ¦  «        z   | _        t          t!          |j                             ¦   «         ¦  «        ¦  «        | _        ddlmŠ ˆfd„| j                             ¦   «         D ¦   «         }|                     d„ ¬¦  «         | xj        t          |¦  «        z  c_        t-          j        | j                             d¦  «        ¦  «                             ¦   «         | _        d d d ¦  «         n# 1 swxY w Y   | j        S )	Nz
recursion:)r;   r@   rA   rI   r   re   c                óJ   •— g | ]\  \  }}\  }}t          |‰¦  «        ¯||f‘Œ S rƒ   )r6   )r†   r;   r`   rs   rf   s       €rK   rˆ   z)JITCallable.cache_key.<locals>.<listcomp>  sO   ø€ ð !@ð !@ð !@Ù%8¡Y d¨A±°°aÝ$.¨s°IÑ$>Ô$>ð!@ $¨ ð !@ð !@ð !@rL   c                ó<   — | d         t          | d         ¦  «        fS )Nr   r   )Úrepr)Úitems    rK   r  z'JITCallable.cache_key.<locals>.<lambda>  s   € °T¸!´W½dÀ4ÈÄ7¹m¼mÐ4L€ rL   )r  r,   )rQ  r`  rN  rÞ   Úgetclosurevarsr  rA   r(   ra  rI   r…   ÚparserQ   r[   rK  ÚdictÚsortedrF   r  rÄ   rf   Úsortr<   r=   r>   rO   )rH   rA   Údependencies_finderÚconstexpr_globalsrf   s       @rK   rV   zJITCallable.cache_key  s  ø€ ð Œ_ð 	Nð 	NØŒyÐ$Ø”yð	Nð 	Nð 	Nð 	Nñ 	Nô 	Nð 	Nð 	Nð
 5 T¤]Ð4Ð4ˆDŒIÝÔ.¨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Ô!à6Ð6Ð6Ð6Ð6Ð6ð!@ð !@ð !@ð !@Ø<@Ô<Q×<WÒ<WÑ<YÔ<Yð!@ñ !@ô !@Ðð ×"Ò"Ð'LÐ'LÐ"ÑMÔMÐMØˆIŒI�Ð.Ñ/Ô/Ñ/ˆIŒIÝœ t¤y×'7Ò'7¸Ñ'@Ô'@ÑAÔA×KÒKÑMÔMˆDŒIð'	Nð 	Nð 	Nñ 	Nô 	Nð 	Nð 	Nð 	Nð 	Nð 	Nð 	Nøøøð 	Nð 	Nð 	Nð 	Nð( ŒyÐs   ‰F!¤E1F!Æ!F%Æ(F%c                ó*   — t          | j        ¦  «        S rN   )r`  rV   rP   s    rK   Ú__hash__zJITCallable.__hash__  s   € Ý�D”NÑ#Ô#Ð#rL   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   )r|   rq  r_  r6   ÚModuler.   ÚbodyÚFunctionDef)rH   Útrees     rK   rq  zJITCallable.parse"  sg   € ÝŒy˜œÑ#Ô#ˆÝ˜$¥¤
Ñ+Ô+Ð+Ð+Ð+Ý�4”9‰~Œ~ Ò"Ð"Ð"Ð"Ý˜$œ) Aœ,­¬Ñ8Ô8Ð8Ð8Ð8ØˆrL   c                ó$   — ddl m}  || ¦  «        S )Nr   )Úconstexpr_type)rÄ   r  )rH   r  s     rK   rl   zJITCallable.type)  s$   € à7Ð7Ð7Ð7Ð7Ð7Øˆ~˜dÑ#Ô#Ð#rL   Úhandlesúlist[ir.value]r*   c                ó   — d S rN   rƒ   )rH   r€  s     rK   Ú_flatten_irzJITCallable._flatten_ir.  s   € ØˆrL   c                ó"   — d| _         || _        dS )a"  
        The only method allowed to modify src.
        Bypasses the __setattr__ restriction by calling super().__setattr__ directly.

        Note that it is the callers responsibility to make sure any triton functions that call this function have the `.hash` value reset to None.
        N)r`  r_  )rH   Únew_srcs     rK   Ú_unsafe_update_srczJITCallable._unsafe_update_src1  s   € ð ˆŒ	ØˆŒ	ˆ	ˆ	rL   c                ó    — t          d¦  «        ‚)NzqCannot set attribute 'src' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)ÚAttributeErrorrP   s    rK   Ú_set_srczJITCallable._set_src;  s   € Ýð (ñ )ô )ð 	)rL   c                ó   — | j         S rN   )r_  rP   s    rK   Ú_get_srczJITCallable._get_src@  s
   € ØŒyÐrL   )ÚfgetÚfsetNrï   )r€  r�  r)   r*   )rY   ri   r¹   r:   rj  r»   rV   rx  rq  rl   rƒ  r†  r‰  r‹  rI   rƒ   rL   rK   rU   rU   Ñ  sÝ   € € € € € ð'(ð '(ð '(ðR,ð ,ð ,ð ðð ð ñ „Xðð0$ð $ð $ðð ð ð ð$ð $ñ „Xð$ðð ð ð ðð ð ð)ð )ð )ð
ð ð ð ˆ(˜ xÐ
0Ñ
0Ô
0€C€C€CrL   rU   c                  ó.   — e Zd ZU ded<   ded<   ded<   dS )ÚJitFunctionInfor   r
  r[   r;   r  r  N)rY   ri   r¹   rm   rƒ   rL   rK   r�  r�  I  s3   € € € € € € àÐÐÑØ€I€I�IØÐÐÑÐÐrL   r�  c                óÜ   ‡— t          |¦  «        t          |¦  «        f}|                      |d ¦  «        }|�|S ˆfd„Št           ‰|¦  «        ¦  «        t          |¦  «        z   }|| |<   |S )Nc                ó4  •— t          | t          ¦  «        rˆfd„| D ¦   «         S t          | ¦  «        rˆfd„| D ¦   «         } | j        |Ž S t          | t          ¦  «        rt	          ˆfd„| D ¦   «         ¦  «        S t          | t
          ¦  «        r| j        S | S )Nc                ó&   •— g | ]} ‰|¦  «        ‘ŒS rƒ   rƒ   ©r†   r•   Úreplace_callabless     €rK   rˆ   z@compute_cache_key.<locals>.replace_callables.<locals>.<listcomp>Y  s%   ø€ Ð:Ð:Ð:¨sÐ%Ð% cÑ*Ô*Ð:Ð:Ð:rL   c                ó&   •— g | ]} ‰|¦  «        ‘ŒS rƒ   rƒ   r“  s     €rK   rˆ   z@compute_cache_key.<locals>.replace_callables.<locals>.<listcomp>[  s%   ø€ Ð=Ð=Ð=°#Ð(Ð(¨Ñ-Ô-Ð=Ð=Ð=rL   c              3  ó.   •K  — | ]} ‰|¦  «        V — Œd S rN   rƒ   r“  s     €rK   ú	<genexpr>z?compute_cache_key.<locals>.replace_callables.<locals>.<genexpr>^  s/   øè è € Ð?Ð?°CÐ*Ð*¨3Ñ/Ô/Ð?Ð?Ð?Ð?Ð?Ð?rL   )r6   r1   r   rJ   ÚtuplerU   rV   )r   Úresultsr”  s     €rK   r”  z,compute_cache_key.<locals>.replace_callablesW  s»   ø€ Ý�c�4Ñ Ô ð 	!Ø:Ð:Ð:Ð:°cÐ:Ñ:Ô:Ð:Ý˜3ÑÔð 	!Ø=Ð=Ð=Ð=¸Ð=Ñ=Ô=ˆGØ �3”= 'Ð*Ð*Ý˜�UÑ#Ô#ð 	!ÝÐ?Ð?Ð?Ð?¸3Ð?Ñ?Ô?Ñ?Ô?Ð?Ý˜�[Ñ)Ô)ð 	!Ø”=Ð Øˆ
rL   )r˜  r[   ry   )Úkernel_key_cacher>  r  r  rV   r”  s        @rK   Úcompute_cache_keyr›  P  sŒ   ø€ Ý�Ñ Ô ¥# g¡,¤,Ð
/€CØ ×$Ò$ S¨$Ñ/Ô/€IØÐØÐð
ð 
ð 
ð 
ð 
õ Ð%Ð% nÑ5Ô5Ñ6Ô6½¸W¹¼ÑE€IØ%Ð�SÑØÐrL   c                óœ   — t          | t          ¦  «        s| S t          | ¦  «        D ]\  }}t          |¦  «        | |<   Œt	          | ¦  «        S rN   )r6   r1   Ú	enumerateÚconvert_to_tuple_if_listr˜  )ro  ÚiÚnested_values      rK   rž  rž  h  sS   € å�d�DÑ!Ô!ð Øˆõ % T™?œ?ð 9ð 9‰ˆˆ<Ý*¨<Ñ8Ô8ˆˆQ‰ˆå�‰;Œ;ÐrL   c                  ó¬   ‡ — e Zd Zd„ Zd"d„Zd„ Zd„ Zd„ Zd„ Zd	„ Z		 	 d#ˆ fd„	Z
d„ Zd„ Zd$d„Zered%d&d„¦   «         Ze	 d%d'd„¦   «         Zd%d „Zd!„ Zˆ xZS )(r  c                ó   — dS )NFrƒ   rP   s    rK   Úis_gluonzJITFunction.is_gluonv  s   € ØˆurL   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 )Nr'  c                ó,   — g | ]\  }}|j         › d |› �‘ŒS )z: ©r;   )r†   rÔ   rÐ   s      rK   rˆ   z*JITFunction._call_hook.<locals>.<listcomp>Š  s,   € Ð_Ð_Ð_¹¸À %¤*Ð4Ð4°Ð4Ð4Ð_Ð_Ð_rL   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  rn  r  ÚcompileÚis_manual_warmupÚalready_compiled)r  r¹   ri   r:  r8  Úparamsrª  r«  r¬  r­  r®  rF  r"  r¯  r�  )rH   Úhookr  r  r­   r©  r  r  r°  r²  r;   r
  Ú	arg_reprsrn  Ú	full_namer±  rþ   s                    rK   Ú
_call_hookzJITFunction._call_hooky  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Ø<BñDô DÐð #ØØ"Ø Ô*ØÔ(Ø!Ô,Ø 'Ô 8Ø'.Ô'FØ"Ô.ØØ#6Ø"ð
ð 
ˆð ˆtØØÝ˜v t¨TÑ2Ô2Ø˜CÐ* 6Ð*Ø&Ø"ð
ñ 
ô 
ð 	
rL   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)rp   Úpre_run_hooksr9  )rH   r·  s     rK   Úadd_pre_run_hookzJITFunction.add_pre_run_hook©  s3   € õ
 ˜‰~Œ~ÐÐÐØÔ×!Ò! $Ñ'Ô'Ð'Ð'Ð'rL   c                óÚ   — ddl m}m}m}m} t
          j                             ¦   «         } ||¦  «        }|| _        || _        || _        t          | j	        | j
        |¦  «        }i i |||fS )z1
        Precompute as much as possible.
        r   )ÚCompiledKernelr³  Ú	ASTSourceÚmake_backend)Úcompilerr¿  r³  rÀ  rÁ  r   ÚactiveÚget_current_targetrB  r  r¶  )rH   r¿  r³  rÀ  rÁ  r­   r5  Úbinders           rK   Ú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ˆØ�2�v˜w¨Ð.Ð.rL   c                óÞ  ‡‡‡— ‰                      |¦  «        }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 ¦   «         }|||
|fS )Nc                ó   — g | ]	}|j         ‘Œ
S rƒ   r§  r  s     rK   rˆ   z*JITFunction._pack_args.<locals>.<listcomp>Â  s   € Ð/Ð/Ð/˜a�1”6Ð/Ð/Ð/rL   c                ó   — g | ]
}|d          ‘ŒS )r   rƒ   r  s     rK   rˆ   z*JITFunction._pack_args.<locals>.<listcomp>Ã  s   € Ð0Ð0Ð0˜A�1�Q”4Ð0Ð0Ð0rL   c                ó   — i | ]\  }}||“Œ	S rƒ   rƒ   )r†   r^   Úvs      rK   r  z*JITFunction._pack_args.<locals>.<dictcomp>Ä  s   € Ð>Ð>Ð>™f˜q !�Q˜Ð>Ð>Ð>rL   Ú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ƒ   )r`   rs   s     rK   r  z(JITFunction._pack_args.<locals>.<lambda>Í  s   € ¸3À+Ò;M€ rL   c           	     óp   •— i | ]2}|t          t          ‰                     ¦   «         ¦  «        |¦  «        “Œ3S rƒ   )r   r1   rã   )r†   ÚpathÚ
bound_argss     €rK   r  z*JITFunction._pack_args.<locals>.<dictcomp>Î  s;   ø€ ÐfÐfÐfÐSW�dÕ-­d°:×3DÒ3DÑ3FÔ3FÑ.GÔ.GÈÑNÔNÐfÐfÐfrL   c                ó:   — g | ]}|d          dk    rdn|d         ‘ŒS )r   rf   rj   r   rƒ   r  s     rK   rˆ   z*JITFunction._pack_args.<locals>.<listcomp>Ð  s.   € ÐPÐPÐP¸A˜!˜Aœ$ +Ò-Ð-�B�B°1°Q´4ÐPÐPÐPrL   c                ó,   — t          |t          ¦  «        S rN   )r6   r[   )r`   r  s     rK   r  z(JITFunction._pack_args.<locals>.<lambda>Ñ  s   € µZÀÅ3Ñ5GÔ5G€ rL   c           	     óX   •— i | ]&}|‰                      t          ‰|¦  «        ¦  «        “Œ'S rƒ   )Ú
parse_attrr   )r†   r^   Úattrvalsr5  s     €€rK   r  z*JITFunction._pack_args.<locals>.<dictcomp>Ò  s4   ø€ ÐVÐVÐVÈ1��G×&Ò&Õ'8¸À1Ñ'EÔ'EÑFÔFÐVÐVÐVrL   )Úparse_optionsr¶  r8  r  ÚKeyErrorr   )rH   r5  rþ   rÑ  r>  r  ÚsigkeysÚsigvalsr  r^   Ú
constexprsr  rÖ  s    ` `        @rK   Ú
_pack_argszJITFunction._pack_args¾  si  øøø€ à×'Ò'¨Ñ/Ô/ˆà/Ð/ 4¤;Ð/Ñ/Ô/ˆØ0Ð0 Ð0Ñ0Ô0ˆØ>Ð>­¨G°WÑ(=Ô(=Ð>Ñ>Ô>ˆ	à FÐ*Ð*Ð*Ð,kÑ*Ô*Ð*Ø˜vÐ%Ð%Ð%Ð'aÑ%Ô%Ð%Ø˜vÐ%Ð%Ð%Ð'aÑ%Ô%Ð%Øð 	Yð 	YˆAØ˜Ô(Ð(Ð(¨Q°gÐ-=Ð-=ÝÐSÐVWÑWÑXÔXÐXøå" 7Ð,MÐ,MÑNÔNˆ
ØfÐfÐfÐfÐ[eÐfÑfÔfˆ
àPÐPÀÐPÑPÔPˆÝ˜hÐ(GÐ(GÑHÔHˆØVÐVÐVÐVÐVÐPUÐVÑVÔVˆà˜	 :¨uÐ4Ð4rL   c               ó(  — |                      d| j        ¦  «        pt          j        j        |d<   t          j        j        |d<   t          j                             ¦   «         }t          j         	                    |¦  «        }| j
        D ]
} ||i |¤Ž Œ| j        |         \  }}	}
}} ||i |¤Ž\  }}}t          j        j        �:t          j                             ¦   «         \  }}|                     d|› d�¦  «         t          |	||¦  «        }|                      |d ¦  «        }|€=|                      |||||¦  «        \  }}}}|                      |||||||¦  «        }|€d S t#          ¦   «         }| j                             ¦   «         D ]?\  \  }}\  }}|                      ||¦  «        x}|k    rt)          d|› d|› d|› �¦  «        ‚Œ@|sÆ|€J ‚t+          |¦  «        r ||¦  «        }t-          |¦  «        }|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ÚdebugÚinstrumentation_modez("custom_pipeline", r$  rS   z1 has changed since we compiled this kernel, from z to r   r   r   )ry   rÞ  r   ÚruntimeÚcompilationrß  r   rÃ  Úget_current_deviceÚget_current_streamr¼  Údevice_cachesÚadd_stages_inspection_hookr9  r›  rÜ  Ú_do_compileÚobjectrF   r  rX   rp   r.   Úlaunch_metadatarã   rö   ÚfunctionÚpacked_metadataÚlaunch_enter_hookÚlaunch_exit_hook)"rH   rù   rú   r—   rþ   r©  rÍ  r·  Úkernel_cacherš  r­   r5  rÅ  rÑ  r>  r  Úinspect_stages_keyÚinspect_stages_hashr  Úkernelr  rÛ  r  Únot_presentr;   r`   rs   Úglobals_dictÚnewValÚ	grid_sizeÚgrid_0Úgrid_1Úgrid_2rè  s"                                     rK   rö   zJITFunction.runÖ  s  € Ø Ÿ*š* W¨d¬jÑ9Ô9ÐP½U¼]Ô=Pˆˆw‰Ý).Ô):Ô)OˆÐ%Ñ&õ ”×1Ò1Ñ3Ô3ˆÝ”×1Ò1°&Ñ9Ô9ˆð Ô&ð 	"ð 	"ˆDØˆD�$Ð!˜&Ð!Ð!Ð!Ð!àBFÔBTÐU[ÔB\Ñ?ˆÐ&¨°¸ð /5¨f°dÐ.E¸fÐ.EÐ.EÑ+ˆ
�N Gõ Œ=Ô3Ð?Ý6;´m×6^Ò6^Ñ6`Ô6`Ñ3ÐÐ 3Ø×!Ò!Ð"OÐ9LÐ"OÐ"OÐ"OÑPÔPÐPåÐ 0°.À'ÑJÔJˆØ×!Ò! # tÑ,Ô,ˆð ˆ>Ø48·O²OÀGÈVÐU_ÐaoØDKñ5Mô 5MÑ1ˆG�Y 
¨Eð ×%Ò% c¨9°f¸jÈ'ÐSXÐZ`ÑaÔaˆFØˆ~Ø�tõ ‘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àˆrL   c                óH   — | j         €| j        n|                       |¦  «        S rN   )Ú_reprrN  )rH   r`   s     rK   rn  zJITFunction.repr  s   € Ø $¤
Ð 2ˆtŒ}ˆ}¸¿
º
À1¹¼ÐErL   Nc	           	     ó¶  •— |r|ng }|r|ng }t          ¦   «                              |¦  «         |j        | _        || _        || _        || _        || _        || _        | t          | j        › d| j
        j        › �<   g | _        t          | j        j                             ¦   «         ¦  «        D ]I\  }	}
|	|v p|
j        |v }|	|v p|
j        |v }| j                             t'          |	|
||¦  «        ¦  «         ŒJt)          | j        ¦  «        | _        d | _        || _        || _        d„ | j        D ¦   «         | _        d„ | j        D ¦   «         | _        g | _        d S )Nr	  c                ó   — g | ]	}|j         ‘Œ
S rƒ   r§  ©r†   Úps     rK   rˆ   z(JITFunction.__init__.<locals>.<listcomp>2  s   € Ð6Ð6Ð6 Q˜!œ&Ð6Ð6Ð6rL   c                ó*   — g | ]}|j         ¯	|j        ‘ŒS rƒ   )rè   rÓ   rü  s     rK   rˆ   z(JITFunction.__init__.<locals>.<listcomp>3  s!   € ÐHÐHÐH Q¸¼ÐH˜1œ5ÐHÐHÐHrL   )r9   r:   ri   r
  ÚversionrÖ   rØ   rù  rè  Ú_triton_jit_function_registryr  r¹   r¶  r�  r  r7  rã   r;   r9  rÒ   r   rÆ  rä  rð  rÞ  rT   Ú	arg_namesrÛ  r¼  )rH   r  rÿ  rÖ   rØ   rÞ  rT   rn  rè  rŸ  rÔ   ÚdnsÚdns_oarJ   s                €rK   r:   zJITFunction.__init__  s…  ø€ à1BÐJÐ-Ð-ÈÐØKiÐ)qÐ)GÐ)GÐoqÐ&å‰Œ×Ò˜ÑÔÐØ”mˆŒØˆŒØ!2ˆÔØ.LˆÔ+ØˆŒ
Ø.ˆÔàQUÕ%¨¬Ð&MÐ&M°t´wÔ7KÐ&MÐ&MÑNàˆŒÝ! $¤.Ô";×"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õ )¨Ô);Ñ<Ô<ˆÔð ˆŒØˆŒ
Ø ˆŒð 7Ð6¨$¬+Ð6Ñ6Ô6ˆŒØHÐH¨$¬+ÐHÑHÔHˆŒð  ˆÔÐÐrL   c           	     ób  ‡‡— dd l }dd lmŠ t          j                             ¦   «         }|                     |¦  «        }|d         | j        k    r t          d|d         › d| j        › �¦  «        ‚t          t          |d         ¦  «        }|d         }| j        |         \  }}}}	}|d         }
|j        |
k    rt          d|
› d|› �¦  «        ‚ˆfd„Šˆfd	„t          ||¦  «        D ¦   «         }t          t          |d
         ¦  «        }|d         }t          t          ||¦  «        ¦  «        }d„ |d                              ¦   «         D ¦   «         }d„ |d                              ¦   «         D ¦   «         }|d         }|	                     |¦  «        }|                      ||||||d¬¦  «        S )Nr   r;   zSpecialization data is for z but trying to preload for r  r  r­   c                ól  •— ‰j                              | ¦  «        r‰                      | ¦  «        S t          | t          ¦  «        rad| v r(‰                     t          | d         ¦  «        ¦  «        S d| v r1| d         }|t          v rt          |         S t          d|› d�¦  «        ‚t          | ¦  «        S )Nrf   r  zUnable to resolve JITFunction z for preload)rÎ   Úis_dtyper6   rr  rf   rž  r   rX   )r�   Újf_keyÚtls     €rK   Ú_decode_constantz-JITFunction.preload.<locals>._decode_constantI  s½   ø€ ØŒx× Ò  Ñ'Ô'ð 'Ø—x’x ‘”Ð&Ý˜%¥Ñ&Ô&ð ^Ø %Ð'Ð'ØŸ<š<Õ(@ÀÀ{ÔASÑ(TÔ(TÑUÔUÐUØ! UÐ*Ð*Ø" >Ô2�FØÕ!>Ð>Ð>Ý<¸VÔDÐDÝ&Ð'\ÈÐ'\Ð'\Ð'\Ñ]Ô]Ð]Ý+¨EÑ2Ô2Ð2rL   c                ó.   •— i | ]\  }}| ‰|¦  «        “ŒS rƒ   rƒ   )r†   r  r�   r	  s      €rK   r  z'JITFunction.preload.<locals>.<dictcomp>V  s+   ø€ ÐgÐgÐg±z°s¸E�cÐ+Ð+¨EÑ2Ô2ÐgÐgÐgrL   r  r  c                ó4   — i | ]\  }}|t          |¦  «        “ŒS rƒ   )rž  r  s      rK   r  z'JITFunction.preload.<locals>.<dictcomp>\  s'   € ÐrÐrÐr¹j¸cÀ5�SÕ2°5Ñ9Ô9ÐrÐrÐrrL   r  c                ób   — i | ],\  }}|t          |t          ¦  «        rt          |¦  «        n|“Œ-S rƒ   )r6   r1   r˜  r  s      rK   r  z'JITFunction.preload.<locals>.<dictcomp>]  sG   € ð 
ð 
ð 
á��Uð ¥¨Eµ4Ñ!8Ô!8ÐC•�u‘”�¸eð
ð 
ð 
rL   r  r  T)rú   )r  Útriton.languagerÅ   r   rÃ  râ  ÚloadsrN  rX   rû   r˜  rä  r  r8  rr  r  r×  ræ  )rH   r±  r  r©  Údeserialized_objr  r  r`   r­   r5  Údeserialized_targetrÛ  r  r  r  r  r  r  r	  r  s                     @@rK   ÚpreloadzJITFunction.preload8  s  øø€ ØˆˆˆØ$Ð$Ð$Ð$Ð$Ð$Ý”×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ˆØ#'Ô#5°fÔ#=Ñ ˆˆ1ˆf�g˜qØ.¨xÔ8Ðð Œ?Ð1Ò1Ð1ÝÐuÐ=PÐuÐuÐmsÐuÐuÑvÔvÐvð	3ð 	3ð 	3ð 	3ð 	3ð hÐgÐgÐgÅSÈÐXeÑEfÔEfÐgÑgÔgˆ
Ý�Ð 0°Ô >Ñ?Ô?ˆ
Ø% lÔ3ˆ
Ý•S˜ ZÑ0Ô0Ñ1Ô1ˆð sÐrÐL\Ð]hÔLi×LoÒLoÑLqÔLqÐrÑrÔrˆ	ð
ð 
à.¨yÔ9×?Ò?ÑAÔAð
ñ 
ô 
ˆð ˜uÔ%ˆØ×'Ò'¨Ñ0Ô0ˆØ×ÒØØØØØØØð  ñ 
ô 
ð 	
rL   c                óX  ‡ ‡‡‡‡‡‡‡‡‡‡‡— ‰ j         ‰         \  Š}Š}	}‰                      t          j        j        ‰‰‰‰‰‰‰g‰¦	  «	        rd S ‰                      ‰ ‰‰‰¦  «        Št          j                             ¦   «         }
|
�Tt          ¦   «         Št          ‰|	‰‰¦  «        }ˆˆˆ ˆˆfd„}ˆˆˆˆˆˆˆ ˆˆˆf
d„}|
                     |||¦  «        }|‰‰<   nO‰                      ‰‰‰j        ¬¦  «        }|‰‰<   ‰                      t          j        j        ‰‰‰‰‰‰‰g‰¦	  «	         |S )Nc                 ó@   •— ‰                      ‰‰‰j        ‰ ¬¦  «        S )N)r­   r  Ú	_env_vars)r³  r  )Úenv_varsr  rH   rI   r­   s   €€€€€rK   Úasync_compilez.JITFunction._do_compile.<locals>.async_compile{  s!   ø€ Ø—|’| C°ÀÔ@PÐ\d�|ÑeÔeÐerL   c                ól   •
— | ‰‰<   ‰                      t          j        j        ‰‰‰	‰‰‰‰g‰
¦	  «	         d S rN   )rº  r   rà  Újit_post_compile_hook)rð  r  rÛ  r©  rí  r  r  rH   r  r­   rú   s    €€€€€€€€€€rK   Úfinalize_compilez1JITFunction._do_compile.<locals>.finalize_compile~  sJ   ø€ Ø$*�˜SÑ!Ø—’¥¤Ô CÀSÈ)ÐU[Ð]cÐeoØ '¨%¨°&ñ:ô :ð :ð :ð :rL   )r­   r  )rä  rº  r   rà  Újit_cache_hookrÀ  r   Úactive_modery   r    r   Úsubmitr³  r  r  )rH   r  r  r©  rÛ  r  r  rú   r`   r5  Ú
async_moderV   r  r  rð  r  rí  rI   r­   s   ````````       @@@@rK   ræ  zJITFunction._do_compilem  s¯  øøøøøøøøøøøø€ Ø.2Ô.@ÀÔ.HÑ+ˆ�a˜ ¨!à�?Š?�5œ=Ô7¸¸iÈÐQWÐYcÐelÐotÐnuØ!ñ#ô #ð 	à�4Ø�nŠn˜T 9¨j¸%Ñ@Ô@ˆå#Ô/×3Ò3Ñ5Ô5ˆ
ØÐ!å6Ñ8Ô8ˆHÝ% c¨7°G¸XÑFÔFˆIðfð fð fð fð fð fð fð fð fð:ð :ð :ð :ð :ð :ð :ð :ð :ð :ð :ð :ð :ð :ð
  ×&Ò& y°-ÐAQÑRÔRˆFØ &ˆL˜ÑÐà—\’\ #¨f¸gÔ>N�\ÑOÔOˆFØ &ˆL˜ÑØ�OŠO�EœMÔ?ÀÀiÐQWÐY_ÐakÐmtØ"˜G Vñ-ô -ð -àˆrL   rH   ú'JITFunction[Callable[P, R]]'r—   úP.argsrþ   úP.kwargsr%   c                ó    — t          d¦  «        ‚)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rX   ©rH   r—   rþ   s      rK   Ú__call__zJITFunction.__call__Œ  s   € ÝÐWÑXÔXÐXrL   Úinstancer*   ÚownerúOptional[type]ú'JITFunction[T]'c                ó   — d S rN   rƒ   ©rH   r$  r%  s      rK   Ú__get__zJITFunction.__get__‘  s   € àˆCrL   ú-'JITFunction[Callable[Concatenate[U, P], R]]'r   úCallable[P, R]c                ó   — d S rN   rƒ   r)  s      rK   r*  zJITFunction.__get__•  s	   € ð ˆCrL   c                ó   — d S rN   rƒ   r)  s      rK   r*  zJITFunction.__get__š  s   € ØˆCrL   c                ó2   — d| j         › d| j        j        › d�S )NzJITFunction(r	  r$  )r
  r  r¹   rP   s    rK   Ú__repr__zJITFunction.__repr__�  s"   € ØC˜dœkÐCÐC¨D¬GÔ,@ÐCÐCÐCÐCrL   )r)   r¤  )NNNNNNN)rH   r  r—   r  rþ   r   r)   r%   rN   )r$  r*   r%  r&  r)   r'  )rH   r+  r$  r   r%  r&  r)   r,  )rY   ri   r¹   r£  rº  r½  rÆ  rÜ  rö   rn  r:   r  ræ  r#  r   r   r*  r0  r¼   r½   s   @rK   r  r  t  sq  ø€ € € € € ðð ð ð.
ð .
ð .
ð .
ð`(ð (ð (ð/ð /ð /ð5ð 5ð 5ð07ð 7ð 7ðrFð Fð Fð mqØ;?ð$ ð $ ð $ ð $ ð $ ð $ ðL3
ð 3
ð 3
ðjð ð ð>Yð Yð Yð Yð ð à	ð	ð 	ð 	ð 	ñ 
Œð	ð 
à,0ð	ð 	ð 	ð 	ñ 
Œð	ð	ð 	ð 	ð 	ðDð Dð Dð Dð Dð Dð DrL   r  r  úJITFunction[T]c                ó   — d S rN   rƒ   rE  s    rK   Újitr3  ¦  s   € à€CrL   ©rÿ  rn  rè  rÖ   rØ   rÞ  rT   rn  úOptional[Callable]rè  rÖ   úOptional[Iterable[int | str]]rØ   rÞ  úOptional[bool]rT   úCallable[[T], JITFunction[T]]c                ó   — d S rN   rƒ   r4  s          rK   r3  r3  «  s	   € ð €CrL   úOptional[T]úKernelInterface[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
    r  r#   r)   r1  c           
     ó¨   •— t          | ¦  «        sJ ‚t          j        j        rddlm}  || ‰‰‰‰‰‰‰¬¦  «        S t          | ‰‰‰‰‰‰‰¬¦  «        S )Nr   )ÚInterpretedFunction)rÿ  rÖ   rØ   rÞ  rT   rn  rè  )rp   r   rà  Ú	interpretÚinterpreterr>  r  )	r  r>  rÞ  rÖ   rØ   rè  rT   rn  rÿ  s	     €€€€€€€rK   Ú	decoratorzjit.<locals>.decoratorÖ  s˜   ø€ Ý˜‰|Œ|ÐÐÐÝŒ=Ô"ð 	Ø8Ð8Ð8Ð8Ð8Ð8Ø&Ð& r°7ÐN_ØFdÐlqØ08¸tÐUdðfñ fô fð fõ ØØØ"3Ø/MØØ!ØØ /ð	ñ 	ô 	ð 	rL   N©r  r#   r)   r1  rƒ   )	r  rÿ  rn  rè  rÖ   rØ   rÞ  rT   rA  s	    ``````` rK   r3  r3  ¹  sb   øøøøøøø€ ð:ð ð ð ð ð ð ð ð ð ð ð ð& 
€~Øˆy˜‰}Œ}Ðð ÐrL   c                  ób   — e Zd ZdZed„ ¦   «         Zd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)rJ   rY   ri   rü   r”   s    rK   rý   zMockTensor.wrap_dtypeû  s/   € àŒ=Ô! WÒ,Ð,°´À7Ò1JÐ1JÝ˜c‘?”?Ð"Øˆ
rL   Nc                ó,   — |€dg}|| _         || _        d S )Nr   )rÎ   Úshape)rH   rÎ   rG  s      rK   r:   zMockTensor.__init__  s    € Øˆ=Ø�CˆEØˆŒ
ØˆŒ
ˆ
ˆ
rL   c                ó    — dg}| j         dd …         D ] }|                     |d         |z  ¦  «         Œ!t          t          |¦  «        ¦  «        S )Nr   rÂ   )rG  r9  r˜  Úreversed)rH   ÚstridesÚsizes      rK   ÚstridezMockTensor.stride  sV   € Ø�#ˆØ”J˜q˜r˜r”Nð 	/ð 	/ˆDØ�NŠN˜7 2œ;¨Ñ-Ñ.Ô.Ð.Ð.Ý•X˜gÑ&Ô&Ñ'Ô'Ð'rL   c                 ó   — dS ©Nr   rƒ   rƒ   rL   rK   Údata_ptrzMockTensor.data_ptr  ó   € àˆqrL   c                 ó   — dS rN  rƒ   rƒ   rL   rK   Ú	ptr_rangezMockTensor.ptr_range  rP  rL   rN   )
rY   ri   r¹   rº   Ústaticmethodrý   r:   rL  rO  rR  rƒ   rL   rK   rü   rü   õ  s’   € € € € € ðð ð
 ðð ñ „\ðð
ð ð ð ð(ð (ð (ð ðð ñ „\ðð ðð ñ „\ðð ð rL   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 rN   )rÎ   ÚbaseÚdatar©  rG  )rH   rW  rÎ   s      rK   r:   zTensorWrapper.__init__  s1   € ØˆŒ
ØˆŒ	Ø”IˆŒ	Ø”kˆŒØ”Y”_ˆŒ
ˆ
ˆ
rL   c                ó4   — | j                              ¦   «         S rN   )rW  rO  rP   s    rK   rO  zTensorWrapper.data_ptr  s   € ØŒy×!Ò!Ñ#Ô#Ð#rL   c                ó    —  | j         j        |Ž S rN   )rW  rL  )rH   r—   s     rK   rL  zTensorWrapper.stride"  s   € ØˆtŒyÔ Ð&Ð&rL   r)   r[   c                ó(   — d| j         › d| j        › d�S )NzTensorWrapper[r¨  r$  )rÎ   rW  rP   s    rK   Ú__str__zTensorWrapper.__str__%  s   € Ø: ¤
Ð:Ð:¨d¬iÐ:Ð:Ð:Ð:rL   c                ó4   — | j                              ¦   «         S rN   )rW  Úelement_sizerP   s    rK   r^  zTensorWrapper.element_size(  s   € ØŒy×%Ò%Ñ'Ô'Ð'rL   c                óZ   — t          | j                             ¦   «         | j        ¦  «        S rN   )rU  rW  ÚcpurÎ   rP   s    rK   r`  zTensorWrapper.cpu+  s   € Ý˜TœYŸ]š]™_œ_¨d¬jÑ9Ô9Ð9rL   c                óD   — | j                              |j         ¦  «         d S rN   )rW  Úcopy_)rH   Úothers     rK   rb  zTensorWrapper.copy_.  s   € ØŒ	�Š˜œ
Ñ#Ô#Ð#Ð#Ð#rL   c                óZ   — t          | j                             ¦   «         | j        ¦  «        S rN   )rU  rW  ÚclonerÎ   rP   s    rK   re  zTensorWrapper.clone1  s    € Ý˜TœYŸ_š_Ñ.Ô.°´
Ñ;Ô;Ð;rL   c                ó\   — t          | j                             |¦  «        | j        ¦  «        S rN   )rU  rW  ÚtorÎ   )rH   r©  s     rK   rg  zTensorWrapper.to4  s"   € Ý˜TœYŸ\š\¨&Ñ1Ô1°4´:Ñ>Ô>Ð>rL   c                ó\   — t          | j                             |¦  «        | j        ¦  «        S rN   )rU  rW  Ú	new_emptyrÎ   )rH   Úsizess     rK   ri  zTensorWrapper.new_empty7  s$   € Ý˜TœY×0Ò0°Ñ7Ô7¸¼ÑDÔDÐDrL   Nrï   )rY   ri   r¹   r:   rO  rL  r\  r^  r`  rb  re  rg  ri  rƒ   rL   rK   rU  rU    s´   € € € € € ð%ð %ð %ð$ð $ð $ð'ð 'ð 'ð;ð ;ð ;ð ;ð(ð (ð (ð:ð :ð :ð$ð $ð $ð<ð <ð <ð?ð ?ð ?ðEð Eð Eð Eð ErL   rU  c                ó  — t          | t          ¦  «        r,|| j        j        k    r| j        S t          | j        |¦  «        S t	          | d¦  «        rt          | |¦  «        S t          dt          | ¦  «        › d�¦  «        ‚)NrO  zCannot reinterpret a rD  )r6   rU  rW  rÎ   Úhasattrr±   rl   )ÚtensorrÎ   s     rK   Úreinterpretrn  ;  s„   € Ý�&�-Ñ(Ô(ð AØ�F”KÔ%Ò%Ð%à”;Ðõ ! ¤¨eÑ4Ô4Ð4Ý	�˜Ñ	$Ô	$ð Aå˜V UÑ+Ô+Ð+åÐ?µ°V±´Ð?Ð?Ð?Ñ@Ô@Ð@rL   c                ó�   — |}t          | ¦  «        D ]3\  }}|                     ¦   «                              d¦  «        r||z  } nŒ4|S )Nzdef )r�  rÇ   rÈ   )rJ  rK  rV  ÚidxÚlines        rK   rU  rU  J  s`   € Ø/Ðõ ˜wÑ'Ô'ð ð ‰	ˆˆTØ�:Š:‰<Œ<×"Ò" 6Ñ*Ô*ð 	Ø  CÑ'Ð ØˆEð	ð  ÐrL   c                óª   — t                                | ¦  «        }|st          d¦  «        ‚t          |                     d¦  «        ¦  «        }|dz  }|S )Nz'No function definition found for kernelÚindentr   )ÚINDENT_PATTERNr\  rM  r.   Úgroup)rc  Úindented_defrX  s      rK   rW  rW  X  sZ   € å!×(Ò(¨Ñ5Ô5€LØð DÝÐBÑCÔCÐCå˜l×0Ò0°Ñ:Ô:Ñ;Ô;Ðà˜1ÑÐØÐrL   c                  ó0   — e Zd Zd„ Zed„ ¦   «         Zd„ ZdS )ÚBoundConstexprFunctionc                ó"   — || _         || _        d S rN   )Ú__self__Ú__func__)rH   r$  r  s      rK   r:   zBoundConstexprFunction.__init__f  s   € Ø ˆŒØˆŒˆˆrL   c                ó   — | j         j        S rN   )r{  rV   rP   s    rK   rV   z BoundConstexprFunction.cache_keyj  s   € àŒ}Ô&Ð&rL   c                ó.   —  | j         | j        g|¢R i |¤ŽS rN   )r{  rz  r"  s      rK   r#  zBoundConstexprFunction.__call__n  s&   € ØˆtŒ}˜Tœ]Ð<¨TÐ<Ð<Ð<°VÐ<Ð<Ð<rL   N)rY   ri   r¹   r:   r»   rV   r#  rƒ   rL   rK   rx  rx  d  sM   € € € € € ðð ð ð ð'ð 'ñ „Xð'ð=ð =ð =ð =ð =rL   rx  c                  óH   ‡ — e Zd Zˆ fd„Zd„ Zedd„¦   «         Zddœd„Zˆ xZS )ÚConstexprFunctionc                óJ   •— t          ¦   «                              |¦  «         d S rN   )r9   r:   )rH   r  rJ   s     €rK   r:   zConstexprFunction.__init__t  s!   ø€ Ý‰Œ×Ò˜ÑÔÐÐÐrL   c                ó*   — |�t          || ¦  «        S | S rN   )rx  )rH   r   Úobjclasss      rK   r*  zConstexprFunction.__get__w  s   € àˆ?Ý)¨#¨tÑ4Ô4Ð4ØˆrL   rH   ú#'ConstexprFunction[Callable[P, R]]'r—   r  rþ   r   r)   r%   c                ó   — d S rN   rƒ   r"  s      rK   r#  zConstexprFunction.__call__}  s   € àˆrL   N)Ú	_semanticc               óÎ   ‡— ddl mŠm} ˆfd„|D ¦   «         }ˆfd„|                     ¦   «         D ¦   «         } | j        |i |¤Ž}|€|S t
          j        j        r|S  ||¦  «        S )Nr   )Ú_unwrap_if_constexprrf   c                ó&   •— g | ]} ‰|¦  «        ‘ŒS rƒ   rƒ   )r†   r  r‡  s     €rK   rˆ   z.ConstexprFunction.__call__.<locals>.<listcomp>„  s%   ø€ Ð6Ð6Ð6¨AÐ$Ð$ QÑ'Ô'Ð6Ð6Ð6rL   c                ó.   •— i | ]\  }}| ‰|¦  «        “ŒS rƒ   rƒ   )r†   r^   rË  r‡  s      €rK   r  z.ConstexprFunction.__call__.<locals>.<dictcomp>…  s+   ø€ ÐJÐJÐJ±°!°Q�!Ð)Ð)¨!Ñ,Ô,ÐJÐJÐJrL   )rÄ   r‡  rf   r  r  r   rà  r?  )rH   r…  r—   rþ   rf   Úresr‡  s         @rK   r#  zConstexprFunction.__call__�  sš   ø€ ØHÐHÐHÐHÐHÐHÐHÐHà6Ð6Ð6Ð6°Ð6Ñ6Ô6ˆØJÐJÐJÐJ¸6¿<º<¹>¼>ÐJÑJÔJˆð ˆdŒg�tÐ&˜vÐ&Ð&ˆàÐàˆJõ Œ=Ô"ð 	ØˆJØˆy˜‰~Œ~ÐrL   )rH   rƒ  r—   r  rþ   r   r)   r%   )rY   ri   r¹   r:   r*  r   r#  r¼   r½   s   @rK   r  r  r  s…   ø€ € € € € ðð ð ð ð ðð ð ð ðð ð ñ „Xðð )-ð ð ð ð ð ð ð ð ð rL   r  úConstexprFunction[T]c                ó    — t          | ¦  «        S )z©
    Wraps an arbitrary Python function so that it can be called at
    compile-time on constexpr arguments in a Triton function and
    returns a constexpr result.
    )r  rE  s    rK   Úconstexpr_functionr�  ”  s   € õ ˜RÑ Ô Ð rL   rï   )FrB  )rn  r5  rè  r5  rÖ   r6  rØ   r6  rÞ  r7  rT   r7  r)   r8  rN   )r  r:  rn  r5  rè  r5  rÖ   r6  rØ   r6  rÞ  r7  rT   r7  r)   r;  )r  r#   r)   r‹  )RÚ
__future__r   r   r|   r7   r<   rÞ   rŸ   rO  r[  rY  Úcollectionsr   Údataclassesr   Ú	functoolsr   Útypingr   r	   r
   r   r   r   r   r   r   r   r   r   Útriton.backendsr   Útypesr   rj   r   r   r   Ú_utilsr   r   r   r   Úcacher   Útriton._C.libtritonr    r!   r"   rD   rC   r³  r]  rt  r#   r$   r%   r&   ÚNodeVisitorr(   rÊ   rÒ   ró   rõ   r"  rB  rF  rU   r   r�  r›  rž  r  r3  rü   rU  rn  rU  rW  rx  r  r�  rƒ   rL   rK   ú<module>r™     s®  ðØ ,Ð ,Ð ,Ð ,Ð ,Ð ,Ð ,Ð ,Ø 
€
€
€
Ø €€€Ø €€€Ø €€€Ø Ð Ð Ð Ø Ð Ð Ð Ø 	€	€	€	Ø €€€Ø #Ð #Ð #Ð #Ð #Ð #Ø !Ð !Ð !Ð !Ð !Ð !Ø %Ð %Ð %Ð %Ð %Ð %ð Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eð  Eà 'Ð 'Ð 'Ð 'Ð 'Ð 'Ø Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð Ø `Ð `Ð `Ð `Ð `Ð `Ð `Ð `Ð `Ð `Ð `Ð `Ø  Ð  Ð  Ð  Ð  Ð  Ø [Ð [Ð [Ð [Ð [Ð [Ð [Ð [Ð [Ð [à!€Ø3€à�”Ð@À"Ä,ÑOÔO€à€GˆC�L„L€Ø€Iˆc�N„N€Ø€GˆC�L„L€Ø€GˆC�L„L€ðf!ð f!ð f!ð f!ð f!˜œñ f!ô f!ð f!ð\Dð Dð Dð Dð4/>ð />ð />ð />ð />ñ />ô />ð />ðdTð Tð Tð TðZð Zð Zð Zð Z�g˜a”jñ Zô Zð Zð&ð ð ð$>*ð >*ð >*ðB0ð 0ð 0ðr1ð r1ð r1ð r1ð r1ñ r1ô r1ð r1ðj !#Ð ð ðð ð ð ð ñ ô ñ „ððð ð ð0	ð 	ð 	ðjDð jDð jDð jDð jD�+˜¨qÔ1ñ jDô jDð jDðd	 
ðð ð ñ 
„ðð 
ð Ø#Ø*.Ø7;ØDHØ Ø#ð
ð 
ð 
ð 
ð 
ñ 
„ð
ð ð4ð Ø#Ø*.Ø7;ØDHØ Ø#ð4ð 4ð 4ð 4ð 4ð 4ðxð ð ð ð ñ ô ð ðB"Eð "Eð "Eð "Eð "Eñ "Eô "Eð "EðJAð Að Að ð  ð  ð	ð 	ð 	ð=ð =ð =ð =ð =˜[ñ =ô =ð =ðð ð ð ð ˜ W¨Q¤Zñ ô ð ðD!ð !ð !ð !ð !ð !rL   