Ë
    Ûéñi|   ã            	      óö  — 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 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, dZ-dZ. ed«      Z/ G d„ dej`                  «      Z1d6d„Z2 G d„ d«      Z3d7d„Z4 G d„ dee/   «      Z5d„ Z6d„ Z7d„ Z8 G d„ d «      Z9e G d!„ d"«      «       Z:d#„ Z;d$„ Z< G d%„ d&e9e5e/   «      Z=ed8d'„«       Z>edddddddd(œ	 	 	 	 	 	 	 	 	 	 	 	 	 d9d)„«       Z>	 d:dddddddd(œ	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d;d*„Z> G d+„ d,«      Z? G d-„ d.«      Z@d/„ ZAd0„ ZB G d1„ d2e9«      ZC G d3„ d4e9«      ZDd5„ ZEy)<é    )ÚannotationsÚdivisionN)Údefaultdict)Ú	dataclass)Úcached_property)	ÚCallableÚGenericÚIterableÚOptionalÚ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_implztriton.languagez"triton.experimental.gluon.languageÚTc                  ó~   ‡ — e Zd ZdZdˆ fd„Zed„ «       Z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).
    c                óò   •— t         ‰| �  «        || _        t        j                  |j                  d«      «      | _        || _        || _        h d£| _	        t        t        ddh| _        i | _        d| _        y )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)Úselfr2   r7   r8   ÚsrcÚ	__class__s        €úT/var/www/pod-logistic/pod-ai/venv/lib/python3.12/site-packages/triton/runtime/jit.pyr1   zDependenciesFinder.__init__.   st   ø€ Ü‰ÑÔØˆŒ	Ü—n‘n S§Z¡Z°Ó%8Ó9ˆŒð ˆŒØ"ˆŒò*
ˆÔ&ô ÜØØð	"
ˆÔð" TVˆÔà*/ˆÕ'ó    c                ó6   — | j                   j                  «       S ©N)r6   Ú	hexdigest©r?   s    rB   ÚretzDependenciesFinder.retY   s   € à�{‰{×$Ñ$Ó&Ð&rC   c                ó†   — t        j                  |j                  «      ryt        |dd«      }|j	                  t
        «      S )NTÚ
__module__Ú )ÚinspectÚ	isbuiltinÚfuncr,   Ú
startswithr;   )r?   ÚnoderN   Úmodules       rB   Ú_is_triton_builtinz%DependenciesFinder._is_triton_builtin]   s6   € Ü×Ñ˜TŸY™YÔ'ØÜ˜˜|¨RÓ0ˆØ× Ñ ¤Ó/Ð/rC   c                ó>  — t        |t        «      sJ ‚| j                  j                  «       |j                  j                  «       z  D ]_  }|\  }}| j                  |   \  }}|j                  |   \  }}||k7  sŒ2t	        d|› d|› d| j
                  › d|j                  › d|› d�«      ‚ | j                  j                  |j                  «       |j                  }|t        t        |dd«      «      z  }| j                  j                  |j                  d	«      «       y )
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-   ÚJITCallabler=   ÚkeysÚRuntimeErrorr2   Ú__name__ÚupdateÚ	cache_keyÚstrr,   r6   r5   )r?   rN   ÚkÚvar_nameÚ_Úv1Úv2Úfunc_keys           rB   Ú_update_hashzDependenciesFinder._update_hashc   s/  € Ü˜$¤Ô,Ð,Ð,ð ×&Ñ&×+Ñ+Ó-°×0EÑ0E×0JÑ0JÓ0LÑLò 	ˆAØ‰KˆH�aØ×)Ñ)¨!Ñ,‰EˆB�Ø×)Ñ)¨!Ñ,‰EˆB�Ø�R‹xÜ"Ø& x j°¸B¸4Ð?OÐPT×PYÑPYÈ{ÐZmÐnr×n{Ñn{Ðm|ð  }Tð  UWð  TXð  XOð  Póð ð	ð 	×Ñ×$Ñ$ T×%:Ñ%:Ô;à—>‘>ˆØ”Cœ  j°%Ó8Ó9Ñ9ˆØ�‰×Ñ˜8Ÿ?™?¨7Ó3Õ4rC   c                ó   — ddl m} |�t        |«      t        u ry t	        |dd«      r#|j
                  D ]  }| j                  |«       Œ y t	        |dd«      ry t	        |dd«      dk(  ry t        |t        «      r| j                  |«       y t        |«      r*t        |t        «      st        ||«      st        d	|› �«      ‚| j                  ry |�/t        j                  |«      |f| j                  |t!        |«      f<   y )
Nr   ©Ú	constexprÚ__triton_aggregate__FÚ__triton_builtin__rJ   rK   ztriton.language.extra.libdevicez!Unsupported function referenced: )Úlanguage.corerf   Útyper   r,   Ú
hash_attrsÚrecord_referencer-   rV   rc   ÚcallablerX   r>   r.   Údeepcopyr=   Úid)r?   ÚvalÚvar_dictr2   rf   Úattrs         rB   rl   z#DependenciesFinder.record_referenceu   s÷   € Ý-ð ˆ;œ$˜s›)¤zÑ1Øä�3Ð.°Ô6ØŸ™ò ,�Ø×%Ñ% dÕ+ð,àä�3Ð,¨eÔ4Øô �3˜ bÓ)Ð-NÒNØä�cœ;Ô'Ø×Ñ˜cÔ"Øä�CŒ=¤¨C´Ô!6¼zÈ#ÈyÔ?YÜÐ!BÀ3À%ÐHÓIÐIð ×*Ò*ØàÐÜ;?¿=¹=ÈÓ;MÈxÐ:XˆD×!Ñ! 4¬¨H«Ð"6Ñ7ØrC   c                ó>  ‡ — t        |j                  «      t        j                  u r|j                  S |j                  ‰ j
                  v ry ˆ fd„} ||j                  «      \  }}|j                  ‰ j                  v r|S ‰ j                  |||j                  «       |S )Nc                ó¶   •— ‰j                   j                  | d «      }|�|‰j                   fS ‰j                  j                  | d «      }|�|‰j                  fS y)N©NN)r7   Úgetr8   )r2   rp   r?   s     €rB   Úname_lookupz2DependenciesFinder.visit_Name.<locals>.name_lookup¤   sZ   ø€ Ø—,‘,×"Ñ" 4¨Ó.ˆCØˆØ˜DŸL™LÐ(Ð(Ø—.‘.×$Ñ$ T¨4Ó0ˆCØˆØ˜DŸN™NÐ*Ð*ØrC   )rj   ÚctxÚastÚStorero   Úlocal_namesr9   rl   )r?   rP   rw   rp   rq   s   `    rB   Ú
visit_NamezDependenciesFinder.visit_Nameœ   s€   ø€ Ü�—‘‹>œSŸY™YÑ&Ø—7‘7ˆNà�7‰7�d×&Ñ&Ñ&àô	ñ $ D§G¡GÓ,‰ˆˆXØ�7‰7�d×4Ñ4Ñ4ØˆJà×Ñ˜c 8¨T¯W©WÔ5Øˆ
rC   c                ó^   — |j                   D �cg c]  }| j                  |«      ‘Œ c}S c c}w rE   )ÚeltsÚvisit)r?   rP   Úelts      rB   Úvisit_TuplezDependenciesFinder.visit_Tuple´   s$   € ð ,0¯9©9Ö5 C�—
‘
˜3•Ò5Ð5ùÒ5s   �*c                óf  — | j                  |j                  «      }t        |t        j                  «      r6| j                  |j                  «      }t        |t        j                  «      rŒ6t        |dd«      }|�|| j                  v ry t        ||j                  «      }| j                  |«       |S )NrY   rK   )	r   Úvaluer-   ry   Ú	Attributer,   r<   rr   rl   )r?   rP   ÚlhsÚlhs_namerH   s        rB   Úvisit_Attributez"DependenciesFinder.visit_Attribute¹   sŠ   € Ø�j‰j˜Ÿ™Ó$ˆÜ˜œcŸm™mÔ,Ø—*‘*˜SŸY™YÓ'ˆCô ˜œcŸm™mÕ,ä˜3 
¨BÓ/ˆØˆ;˜( d×&<Ñ&<Ñ<ØÜ�c˜4Ÿ9™9Ó%ˆØ×Ñ˜cÔ"Øˆ
rC   c                ó–   — |j                   j                   D �ch c]  }|j                  ’Œ c}| _        | j                  |«       y c c}w rE   )ÚargsÚargr{   Úgeneric_visit)r?   rP   rŠ   s      rB   Úvisit_FunctionDefz$DependenciesFinder.visit_FunctionDefÄ   s4   € à/3¯y©y¯~©~Ö>¨˜CŸG›GÒ>ˆÔØ×Ñ˜4Õ ùò ?s   ™Ac                óp  ‡ — ˆ fd„}t        j                  |j                  |j                  |j                  r|j                  gng |j
                  «      D ]  }‰ j                  |«       Œ  ||j                  «       |j                  �‰ j                  |j                  «        ||j                  «       y )Nc                ó�   •— 	 ‰j                   rJ ‚d‰_         | D ]  }|€Œ‰j                  |«       Œ 	 d‰_         y # d‰_         w xY w)NTF)r>   r   )ÚdefaultsÚexprr?   s     €rB   Úvisit_defaultsz:DependenciesFinder.visit_arguments.<locals>.visit_defaultsÕ   sS   ø€ ð8Ø×:Ò:Ð:Ð:Ø26�Ô/Ø$ò )�DØÑ'ØŸ
™
 4Õ(ñ)ð 38�Õ/ø°%�Õ/ús   ƒ<  < ¼	A)
Ú	itertoolsÚchainÚposonlyargsr‰   ÚvarargÚ
kwonlyargsr   Úkw_defaultsÚkwargr�   )r?   rP   r‘   rŠ   s   `   rB   Úvisit_argumentsz"DependenciesFinder.visit_argumentsÉ   sŒ   ø€ ô	8ô —?‘? 4×#3Ñ#3°T·Y±YÐQU×Q\ÒQ\ÀÇÁÁÐbdÐfj×fuÑfuÓvò 	ˆCØ�J‰J�s�Oð	ñ 	�t×'Ñ'Ô(à�:‰:Ð!Ø�J‰J�t—z‘zÔ"á�t—}‘}Õ%rC   c                óº   — | j                  |«      }t        |t        «      r| xj                  t	        |«      z  c_        y | j                  j                  |«       y rE   )r   r-   r(   r{   ÚsetÚadd)r?   rP   Útargets      rB   ÚvisitAssnTargetz"DependenciesFinder.visitAssnTargeté   sE   € ð —‘˜DÓ!ˆÜ�fœdÔ#Ø×Ò¤ F£Ñ+Öà×Ñ× Ñ  Õ(rC   c                ó¨   — t        |j                  «      dk7  rt        d«      ‚| j                  |j                  d   «       | j	                  |«       y )Nr   z2Simultaneous multiple assignment is not supported.r   )r%   ÚtargetsÚ	TypeErrorrž   r‹   ©r?   rP   s     rB   Úvisit_AssignzDependenciesFinder.visit_Assignò   sG   € Üˆt�|‰|Ó Ò!ô
 ÐPÓQÐQà×Ñ˜TŸ\™\¨!™_Ô-ð 	×Ñ˜4Õ rC   c                ó\   — | j                  |j                  «       | j                  |«       y rE   ©rž   r�   r‹   r¢   s     rB   Úvisit_AnnAssignz"DependenciesFinder.visit_AnnAssignÿ   ó$   € Ø×Ñ˜TŸ[™[Ô)ð 	×Ñ˜4Õ rC   c                ó\   — | j                  |j                  «       | j                  |«       y rE   r¥   r¢   s     rB   Ú	visit_ForzDependenciesFinder.visit_For  r§   rC   )ÚreturnÚNoneru   )rY   rJ   Ú__qualname__Ú__doc__r1   ÚpropertyrH   rR   rc   rl   r|   r�   r‡   rŒ   r™   rž   r£   r¦   r©   Ú__classcell__©rA   s   @rB   r!   r!   "   s`   ø„ ñ	õ)0ðV ñ'ó ð'ò0ò5ó$%òNò06ò
	ò!ò
&ò@)ò!ò!ö!rC   r!   c                ó   — dd l mc m} t        | t        «      rÇ| j                  «       } | j                  d«      r7| j                  d«      } t        | «      } | j                  d«      sJ ‚d| dd  z   S | j                  d«      rdt        | d d «      z   S | j                  d«      rdt        | dd  «      z   S | j                  d«      r“t        | j                  d«      «      S t        | |j                  «      rdt        | j                  «      › �S t        | |j                  «      r| j                  } n(t        | t        «      r| j                  } nt	        | «      } t!        j"                  | j%                  d	d
«      | «      S )Nr   zconst ÚconstÚ*ú*kr   éÿÿÿÿztl.Ú_trK   )Útriton.language.coreÚlanguageÚcorer-   r\   ÚstriprO   ÚremoveprefixÚ_normalize_tyÚendswithÚpointer_typeÚ
element_tyÚdtyper2   rj   rY   r   rv   Úreplace)Útyr¹   s     rB   r¼   r¼     s?  € ß'Ð'Ü�"”cÔØ�X‰X‹ZˆØ�=‰=˜Ô"Ø—‘ Ó)ˆBÜ˜rÓ"ˆBØ—=‘= Ô%Ð%Ð%Ø˜"˜Q˜R˜&‘=Ð Ø�;‰;�sÔØœ r¨#¨2 wÓ/Ñ/Ð/Ø�=‰=˜ÔØœ r¨!¨" vÓ.Ñ.Ð.Ø�=‰=˜ÔÜ  §¡°Ó!7Ó8Ð8Ü	�B˜×)Ñ)Ô	*Ø”= §¡Ó/Ð0Ð1Ð1Ü	�B˜Ÿ
™
Ô	#Ø�W‰W‰Ü	�BœÔ	Ø�[‰[‰ä�‹WˆÜ%×)Ñ)¨"¯*©*°T¸2Ó*>ÀÓCÐCrC   c                  ó�   — e Zd ZdZ	 	 dd„Zed„ «       Zedd„«       Zedd„«       Zed„ «       Z	ed„ «       Z
ed„ «       Zed	„ «       Zy
)ÚKernelParamzBRepresents a parameter (name plus metadata) to a @jit'ed function.c                ó<   — || _         || _        || _        || _        y rE   )ÚnumÚ_paramÚdo_not_specializeÚdo_not_specialize_on_alignment)r?   rÆ   ÚparamrÈ   rÉ   s        rB   r1   zKernelParam.__init__.  s    € àˆŒØˆŒØ!2ˆÔØ.LˆÕ+rC   c                ó.   — | j                   j                  S rE   )rÇ   r2   rG   s    rB   r2   zKernelParam.name5  s   € à�{‰{×ÑÐrC   c                óÐ   — | j                   j                  r1| j                   j                  t        j                  j                  k(  ryt        | j                   j                  «      S )NrK   )rÇ   Ú
annotationrL   Ú	ParameterÚemptyr¼   rG   s    rB   rÍ   zKernelParam.annotation9  sD   € à�{‰{×%Ò%¨¯©×)?Ñ)?Ä7×CTÑCT×CZÑCZÒ)ZØÜ˜TŸ[™[×3Ñ3Ó4Ð4rC   c                óÌ   — | j                   }|j                  d«      r|dd  }n|j                  d«      r|dd  }|t        t        j                  «       «      v r| j                   S y)Nr´   r   r³   r   rK   )rÍ   rO   r›   r   Úvalues)r?   Úas     rB   Úannotation_typezKernelParam.annotation_type?  s]   € à�O‰OˆØ�<‰<˜ÔØ�!�"�‰AØ�\‰\˜#ÔØ�!�"�ˆAØ”Ô.×5Ñ5Ó7Ó8Ñ8Ø—?‘?Ð"ØrC   c                ó   — d| j                   v S ©Nrf   )rÍ   rG   s    rB   Úis_constexprzKernelParam.is_constexprJ  s   € à˜dŸo™oÐ-Ð-rC   c                ór   — | j                   ryd| j                  v xs | j                  j                  d«      S )NFr²   r´   )rÖ   rÍ   rO   rG   s    rB   Úis_constzKernelParam.is_constN  s1   € à×ÒØØ˜$Ÿ/™/Ð)ÒM¨T¯_©_×-GÑ-GÈÓ-MÐMrC   c                ó.   — | j                   j                  S rE   )rÇ   ÚdefaultrG   s    rB   rÚ   zKernelParam.defaultT  s   € à�{‰{×"Ñ"Ð"rC   c                ód   — | j                   j                  t        j                  j                  k7  S rE   )rÇ   rÚ   rL   rÎ   rÏ   rG   s    rB   Úhas_defaultzKernelParam.has_defaultX  s#   € à�{‰{×"Ñ"¤g×&7Ñ&7×&=Ñ&=Ñ=Ð=rC   N)rÆ   r$   rÊ   zinspect.ParameterrÈ   ÚboolrÉ   rÝ   ©rª   r\   )rY   rJ   r¬   r­   r1   r   r2   rÍ   rÓ   rÖ   rØ   r®   rÚ   rÜ   © rC   rB   rÄ   rÄ   +  s¢   „ ÙLðMØ15óMð ñ ó ð ð ò5ó ð5ð
 òó ðð ñ.ó ð.ð ñNó ðNð
 ñ#ó ð#ð ñ>ó ñ>rC   rÄ   c                ó6   — d}d}t        t        | |||«      d   S )NFTr   )r   r   )rŠ   Ú
specializerØ   Úaligns       rB   Úmangle_typerã   ]  s%   € Ø€HØ€EÜ!¤+¨s°H¸jÈ%ÓPÐQRÑSÐSrC   c                  ó,   — e Zd ZU ded<   d„ Zd„ Zdd„Zy)ÚKernelInterfacer   Úrunc               ó\   —  | j                   t        t        j                  |«      |ddœ|¤ŽS )NT©ÚgridÚwarmup)ræ   ÚmapÚ
MockTensorÚ
wrap_dtype)r?   ré   r‰   Úkwargss       rB   rê   zKernelInterface.warmupf  s*   € Øˆt�x‰x´´Z×5JÑ5JÈDÓ1Q˜T¨$Ñ\ÐU[Ñ\Ð\rC   c               ó   — t        d«      ‚)Nzrun not implemented)ÚNotImplementedError)r?   ré   rê   r‰   rî   s        rB   ræ   zKernelInterface.runi  s   € Ü!Ð"7Ó8Ð8rC   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é   r?   s     €€rB   ú<lambda>z-KernelInterface.__getitem__.<locals>.<lambda>r  s   ø€  x t§x¡xÈ$°TÀ%Ñ'YÐRXÑ'Y€ rC   rß   )r?   ré   s   ``rB   Ú__getitem__zKernelInterface.__getitem__l  s   ù€ ô ZÐYrC   N)rª   r   )rY   rJ   r¬   Ú__annotations__rê   ræ   rô   rß   rC   rB   rå   rå   c  s   … Ø	
ƒFò]ò9ôZrC   rå   c           
     óF  — |j                  «       D ��ci c]S  \  }}||j                  j                  dk(  rt        |«      n(|j                  j                  dk(  rd|j                  in|“ŒU }}}dd l}| ||j                  «       D �cg c]  }t        |«      ‘Œ c}t        |j                  «       «      |j                  «       D �cg c]  }t        |«      ‘Œ c}t        |j                  «       «      |j                  dœ}	|j                  |	«      }
|
S c c}}w c c}w c c}w )NrÀ   rf   r   )r2   Ú	signatureÚconstant_keysÚconstant_valsÚ
attrs_keysÚ
attrs_valsÚoptionsÚkey)ÚitemsrA   rY   r\   rƒ   ÚjsonrW   r(   rÑ   Ú__dict__Údumps)r2   r÷   Ú	constantsÚattrsrü   rý   rƒ   rÿ   ÚxÚobjÚserialized_objs              rB   Úserialize_specialization_datar  v  s   € ð $Ÿ/™/Ó+÷ñ ˆC�ð 	˜5Ÿ?™?×3Ñ3°wÒ>ŒS�ŒZØ&+§o¡o×&>Ñ&>À+Ò&Mˆ�e—k‘kÑ"ÐSXñ	Yð€Iñ ó à 9ÐQZ×Q_ÑQ_ÓQaÖ?bÈAÄÀQÅÒ?bÜˆY×ÑÓÓ À%Ç*Á*Ã,Ö0O¸Q´°aµÒ0OÔ_cÐdi×dpÑdpÓdrÓ_sØ×#Ñ#¨Cñ€Cð
 —Z‘Z “_€NØÐùóùò @cùÚ0Os   ”ADÂDÃDc                ó   — t        | j                  «      t        |«      k(  sJ ‚g }t        | j                  j                  «       |«      D �]  \  }}|j                  r|j                  d|› d�«       Œ)|j                  rdnd}|j                  rdnd}|j                  rdnd}d|› d|› d|› d|› d�	}	|j                  r�t        |j                  t        «      r"|j                  dk(  s|j                  dd	 d
v rd}|r#|j                  d|j                  › d|	› d�«       ŒÕ|j                  d|j                  › d�«       Œõ|j                  |	› «       �Œ	 d„ }
ddj                  t        t        |
| j                  j                  «       «      «      dgz   «      › ddj                  | j                  j                  «       D �cg c]
  }d|› d|› �‘Œ c}«      › ddj                  |«      › d�}| j                  j                  «       D ��ci c];  \  }}|j                   t"        j$                  j&                  urd|› �|j                   “Œ= }}}t(        }||d<   ||d<   t*        |d<   t-        ||«       |d   S c c}w c c}}w )a2  
    Equivalent to sig.bind followed by apply_defaults. This generates a
    native Python function (using exec) which can be memoized on a per-kernel
    basis to avoid having to run these expensive functions -- which constitute
    much of the kernel launch overhead -- every time we run the kernel.
    z("constexpr", ú)ÚTrueÚFalsezspecialize_impl(backend, ú, Úu1Nr   )ÚfpÚbfFz("z",) + z[1:]z", None)c                óx   — | d   j                   t        j                  j                  u r| d   S | d   › d| d   › �S )Nr   r   z	=default_)rÚ   rL   rÎ   rÏ   )r  s    rB   ró   z0create_function_from_signature.<locals>.<lambda>¨  sA   € ˜A˜a™DŸL™L¬G×,=Ñ,=×,CÑ,CÑC�A�a‘D€ ÈAÈaÉDÈ6ÐQZÐ[\Ð]^Ñ[_ÐZ`ÐIa€ rC   z
def dynamic_func(z	**optionsz):
    params = {ú'z': z}
    specialization = [ú,z-]
    return params, specialization, options
Údefault_Úspecialize_implÚbackendrV   Údynamic_func)r%   Ú
parametersÚziprW   rÖ   ÚappendrØ   rÈ   rÉ   rÓ   r-   r\   Újoinr(   rë   rþ   rÚ   rL   rÎ   rÏ   r   rV   Úexec)ÚsigÚkparamsr  Úspecializationr2   ÚkprØ   rá   râ   rH   rŠ   Ú	func_bodyrÊ   Úfunc_namespacer  s                  rB   Úcreate_function_from_signaturer"  ‡  s†  € ô ˆs�~‰~Ó¤# g£,Ò.Ð.Ð.à€Nä˜Ÿ™×+Ñ+Ó-¨wÓ7ó 0‰ˆˆbØ�?Š?Ø×!Ñ! N°4°&¸Ð":Õ;à!#§¢‘v°'ˆHØ$&×$8Ò$8™¸fˆJØ!×@Ò@‘GÀfˆEØ-¨d¨V°2°h°Z¸rÀ*ÀÈRÐPUÈwÐVWÐXˆCØ×!Ò!Ü˜b×0Ñ0´#Ô6Ø×)Ñ)¨TÒ1°R×5GÑ5GÈÈÐ5KÈ|Ñ5[à%*˜
ÙØ"×)Ñ)¨B¨r×/AÑ/AÐ.BÀ&ÈÈÈTÐ*RÕSð #×)Ñ)¨B¨r×/AÑ/AÐ.BÀ(Ð*KÕLà×%Ñ%¨¨Ö/ð'0ñ, b€CðØ—)‘)œD¤ S¨#¯.©.×*>Ñ*>Ó*@Ó!AÓBÀkÀ]ÑRÓSÐTð UØ—	‘	¸3¿>¹>×;NÑ;NÓ;PÖQ°4˜Q˜t˜f C¨ vÒ.ÒQÓRÐSð TØ—x‘x Ó/Ð0ð 1ð€Ið Ÿ>™>×/Ñ/Ó1÷áˆD�%Ø�=‰=¤× 1Ñ 1× 7Ñ 7Ñ7ð �4�&Ð˜5Ÿ=™=Ñ(ð€Nñ ô -€OØ(7€NÐ$Ñ%Ø '€N�9ÑÜ$/€N�=Ñ!ô 	ˆ�NÔ#ð ˜.Ñ)Ð)ùò+ Rùós   Ç I5ÈA I:c                ó8   — | j                   › d| j                  › �S )Nú.)rJ   r¬   ©Úfns    rB   Úget_full_namer'  Ã  s   € Ø�m‰mˆ_˜A˜bŸo™oÐ.Ð/Ð/rC   c                  ól   — e Zd Zd„ Zd„ Zedd„«       Zd„ Zd„ Zed„ «       Z	d„ Z
d„ Zd	„ Z eee¬
«      Zy)rV   c                ó´  — || _         t        j                  |«      | _        	 t        j                  |«      \  | _        | _        t        |«      | _	        t        j                  «       | _        t        j                  dj                  | j                  «      «      }|t!        j"                  d|t         j$                  «      j'                  «       d  }|| _        d | _        i | _        |j.                  | _        |j0                  | _        |j2                  | _        |j4                  | _        |j6                  | _        y # t        $ r}t        d«      |‚d }~ww xY w)Nz1@jit functions should be defined in a Python filerK   z^def\s+\w+\s*\()r&  rL   r÷   ÚgetsourcelinesÚraw_srcÚstarting_line_numberÚOSErrorÚ
ValueErrorr'  Ú_fn_nameÚ	threadingÚRLockÚ
_hash_lockÚtextwrapÚdedentr  ÚreÚsearchÚ	MULTILINEÚstartÚ_srcÚhashr=   r­   rY   r¬   Ú__globals__rJ   )r?   r&  Úer@   s       rB   r1   zJITCallable.__init__É  s  € ØˆŒÜ ×*Ñ*¨2Ó.ˆŒð	YÜ6=×6LÑ6LÈRÓ6PÑ3ˆDŒL˜$Ô3ô & bÓ)ˆŒÜ#Ÿ/™/Ó+ˆŒô �o‰o˜bŸg™g d§l¡lÓ3Ó4ˆØ”"—)‘)Ð.°´R·\±\ÓB×HÑHÓJÐKÐLˆØˆŒ	ØˆŒ	ð TVˆÔð —z‘zˆŒØŸ™ˆŒØŸO™OˆÔØŸ>™>ˆÔØŸ-™-ˆ�øô7 ò 	YÜÐPÓQÐWXÐXûð	Yús   £"D= Ä=	EÅEÅEc                ón   — | j                   t        j                  | j                  «      j                  z  S rE   )r;  rL   Úgetclosurevarsr&  r8   rG   s    rB   Úget_capture_scopezJITCallable.get_capture_scopeë  s(   € Ø×Ñ¤'×"8Ñ"8¸¿¹Ó"A×"KÑ"KÑKÐKrC   c                óÐ  — | j                   5  | j                  �| j                  cd d d «       S d| j                  › �| _        t        j                  | j
                  «      j                  }t        | j                  | j                  || j                  ¬«      }|j                  | j                  «       «       |j                  t        | j                  «      z   | _        t        t!        |j"                  j%                  «       «      «      | _        ddlm} | xj                  t        | j"                  j%                  «       D ���cg c]  \  \  }}\  }}t+        ||«      r||f‘Œ c}}}«      z  c_        t-        j.                  | j                  j1                  d«      «      j3                  «       | _        d d d «       | j                  S c c}}}w # 1 sw Y   | j                  S xY w)Nz
recursion:)r2   r7   r8   r@   r   re   r#   )r2  r:  r/  rL   r>  r&  r8   r!   r;  r@   r   ÚparserH   r\   r,  ÚdictÚsortedr=   rþ   r·   rf   r-   r3   r4   r5   rF   )r?   r8   Údependencies_finderrf   r2   r_   rp   s          rB   r[   zJITCallable.cache_keyî  s„  € ð �_‰_ñ 	NØ�y‰yÐ$Ø—y‘y÷	Nñ 	Nð
 % T§]¡] OÐ4ˆDŒIÜ×.Ñ.¨t¯w©wÓ7×AÑAˆIÜ"4¸$¿-¹-ÐQU×QaÑQaÐmvØ9=¿¹ô#CÐà×%Ñ% d§j¡j£lÔ3Ø+×/Ñ/´#°d×6OÑ6OÓ2PÑPˆDŒIÜ$(¬Ð0C×0TÑ0T×0ZÑ0ZÓ0\Ó)]Ó$^ˆDÔ!å6Ø�IŠIœØ9=×9NÑ9N×9TÑ9TÓ9V÷=ð =Ù"5¡) 4¨©X¨c°1Ü!+¨C°Ô!;ð  $ Sškô =ó >ñ >�Iô  Ÿ™ t§y¡y×'7Ñ'7¸Ó'@ÓA×KÑKÓMˆDŒI÷#	Nð$ �y‰yÐùô	=÷	Nð$ �y‰yÐús$   �G¯DGÅ G
Å%AGÇ
GÇG%c                ó,   — t        | j                  «      S rE   )r:  r[   rG   s    rB   Ú__hash__zJITCallable.__hash__  s   € Ü�D—N‘NÓ#Ð#rC   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   )ry   rA  r9  r-   ÚModuler%   ÚbodyÚFunctionDef)r?   Útrees     rB   rA  zJITCallable.parse  s_   € Ü�y‰y˜Ÿ™Ó#ˆÜ˜$¤§
¡
Ô+Ð+Ð+Ü�4—9‘9‹~ Ò"Ð"Ð"Ü˜$Ÿ)™) A™,¬¯©Ô8Ð8Ð8ØˆrC   c                ó   — ddl m}  || «      S )Nr   )Úconstexpr_type)r·   rM  )r?   rM  s     rB   rj   zJITCallable.type  s   € å7Ù˜dÓ#Ð#rC   c                ó    — d| _         || _        y)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:  r9  )r?   Únew_srcs     rB   Ú_unsafe_update_srczJITCallable._unsafe_update_src  s   € ð ˆŒ	Øˆ�	rC   c                ó   — t        d«      ‚)NzqCannot set attribute 'src' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)ÚAttributeErrorrG   s    rB   Ú_set_srczJITCallable._set_src!  s   € Üð (ó )ð 	)rC   c                ó   — | j                   S rE   )r9  rG   s    rB   Ú_get_srczJITCallable._get_src&  s   € Ø�y‰yÐrC   )ÚfgetÚfsetNrÞ   )rY   rJ   r¬   r1   r?  r®   r[   rF  rA  rj   rP  rS  rU  r@   rß   rC   rB   rV   rV   Ç  s]   „ ò (òDLð òó ðò,$òð ñ$ó ð$òò)ò
ñ ˜ xÔ
0�CrC   rV   c                  ó,   — e Zd ZU ded<   ded<   ded<   y)ÚJitFunctionInfor   rQ   r\   r2   ÚJITFunctionÚjit_functionN)rY   rJ   r¬   rõ   rß   rC   rB   rY  rY  ,  s   … àÓØ
ƒIØÔrC   rY  c                ó®   ‡— t        |«      t        |«      f}| j                  |d «      }|�|S ˆfd„Št         ‰|«      «      t        |«      z   }|| |<   |S )Nc                óF  •— t        | t        «      r| D �cg c]
  } ‰|«      ‘Œ c}S t        | «      r$| D �cg c]
  } ‰|«      ‘Œ }} | j                  |Ž S t        | t        «      rt	        ˆfd„| D «       «      S t        | t
        «      r| j                  S | S c c}w c c}w )Nc              3  ó.   •K  — | ]  } ‰|«      –— Œ y ­wrE   rß   )Ú.0rŠ   Úreplace_callabless     €rB   ú	<genexpr>z?compute_cache_key.<locals>.replace_callables.<locals>.<genexpr>A  s   øè ø€ Ò?°CÑ*¨3×/Ñ?ùs   ƒ)r-   r(   r   rA   ÚtuplerV   r[   )r  rŠ   Úresultsr`  s      €rB   r`  z,compute_cache_key.<locals>.replace_callables:  s’   ø€ Ü�cœ4Ô Ø69Ö:¨sÑ% cÕ*Ò:Ð:Ü˜3ÔØ9<Ö=°#Ñ(¨Õ-Ð=ˆGÐ=Ø �3—=‘= 'Ð*Ð*Ü˜œUÔ#ÜÓ?¸3Ô?Ó?Ð?Ü˜œ[Ô)Ø—=‘=Ð Øˆ
ùò ;ùâ=s
   –B·B)rb  r\   rv   )Úkernel_key_cacher  rü   rý   r[   r`  s        @rB   Úcompute_cache_keyre  3  se   ø€ Ü�Ó ¤# g£,Ð
/€CØ ×$Ñ$ S¨$Ó/€IØÐØÐô
ô Ñ% nÓ5Ó6¼¸W»ÑE€IØ%Ð�SÑØÐrC   c                ó~   — t        | t        «      s| S t        | «      D ]  \  }}t        |«      | |<   Œ t	        | «      S rE   )r-   r(   Ú	enumerateÚconvert_to_tuple_if_listrb  )ÚitemÚiÚnested_values      rB   rh  rh  K  sD   € ä�dœDÔ!Øˆô % T›?ò 9‰ˆˆ<Ü*¨<Ó8ˆˆQŠð9ô �‹;ÐrC   c                  ól   ‡ — 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„ Zd„ Zˆ xZS )rZ  c                 ó   — y)NFrß   rG   s    rB   Úis_gluonzJITFunction.is_gluonY  s   € ØrC   c	                ó²  — |sy | j                   j                  }	| j                   j                  }
dj                  t	        | j
                  |d   «      D ��cg c]  \  }}|j                  › d|› �‘Œ c}}«      }|	› d|j                  › d|j                  › d|j                  › d|j                  › d|j                  › d	|› d
�}t        | j                   «      }t        ||||d   ||«      }||||j                  |j                  |j                  |j                  |j                  |j                  |||dœ} |||t        |
|	| «      d|i|¥|d¬«      S c c}}w )Nr  r   z: z[num_warps=z, num_ctas=z, num_stages=z, enable_fp_fusion=z, launch_cooperative_grid=ú](r	  r   )r÷   Údevicer  Ú	num_warpsÚnum_ctasÚ
num_stagesÚenable_fp_fusionÚlaunch_cooperative_gridÚextern_libsÚconfigsÚspecialization_dataÚ	is_warmuprý   F)rý   Úreprr&  ÚcompileÚis_manual_warmupÚalready_compiled)r&  r¬   rJ   r  r  Úparamsr2   rr  rs  rt  ru  rv  r'  r  rw  rY  )r?   Úhookrý   r÷   rq  r  rü   rx  rz  r2   rQ   rÊ   rÂ   Ú	arg_reprsr{  Ú	full_namery  rî   s                     rB   Ú
_call_hookzJITFunction._call_hook\  s¶  € ñ Øà�w‰w×#Ñ#ˆØ—‘×#Ñ#ˆØ—I‘IÄcÈ$Ï+É+ÐWZÐ[\ÑW]ÓF^×_¹¸À %§*¡* ¨R°¨tÒ4Ó_Ó`ˆ	Ø��{ 7×#4Ñ#4Ð"5°[À×AQÑAQÐ@RÐR_Ð`g×`rÑ`rÐ_sð  tGð  HO÷  H`ñ  H`ð  Gað  a{ð  |C÷  |[ñ  |[ð  {\ð  \^ð  _hð  ^ið  ijð  kˆÜ! $§'¡'Ó*ˆ	ä;¸IÀyÐR[Ð]dÐefÑ]gÐipÐruÓvÐð #ØØ"Ø ×*Ñ*Ø×(Ñ(Ø!×,Ñ,Ø '× 8Ñ 8Ø'.×'FÑ'FØ"×.Ñ.ØØ#6Ø"ñ
ˆñ ØØÜ˜v t¨TÓ2Ø˜CÐ* 6Ð*Ø&Ø"ô
ð 	
ùó+ `s   ÁE
c                óT   — t        |«      sJ ‚| j                  j                  |«       y)z�
        Add a hook that will be executed prior to the execution of run
        function with args and kwargs passed into the kernel
        N)rm   Úpre_run_hooksr  )r?   r€  s     rB   Úadd_pre_run_hookzJITFunction.add_pre_run_hookŠ  s$   € ô
 ˜Œ~Ðˆ~Ø×Ñ×!Ñ! $Õ'rC   c                óà   — ddl m}m}m}m} t
        j                  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_targetr"  r÷   r  )r?   rˆ  r|  r‰  rŠ  r�   r  Úbinders           rB   Úcreate_binderzJITFunction.create_binder’  sd   € ÷ 	PÓOÜ—‘×1Ñ1Ó3ˆÙ˜vÓ&ˆØ,ˆÔØˆŒØ"ˆŒÜ/°·±ÀÇÁÈWÓUˆØ�2�v˜w¨Ð.Ð.rC   c           
     ó´  — |j                  |«      }| j                  D �cg c]  }|j                  ‘Œ }}|D �cg c]  }|d   ‘Œ	 }}t        ||«      D �	�
ci c]  \  }	}
|	|
“Œ
 }}	}
d|vsJ d«       ‚d|vsJ d«       ‚d|vsJ d«       ‚|D ]#  }	|	|j                  vsŒ|	|vsŒt        d|	z  «      ‚ t        |d	„ «      }|D �ci c]&  }|t        t        |j                  «       «      |«      “Œ( }}|D �cg c]  }|d
   ‘Œ	 }}t        |d„ «      }|D �	ci c]  }	|	|j                  t        ||	«      «      “Œ  }}	||||fS c c}w c c}w c c}
}	w c c}w c c}w c c}	w )Nr   Údevice_typez=device_type option is deprecated; current target will be usedrq  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_   rp   s     rB   ró   z(JITFunction._pack_args.<locals>.<lambda>®  s   € ¸3À+Ñ;M€ rC   r   c                ó"   — t        |t        «      S rE   )r-   r\   )r_   r  s     rB   ró   z(JITFunction._pack_args.<locals>.<lambda>²  s   € ´ZÀÄ3Ó5G€ rC   )Úparse_optionsr  r2   r  r   ÚKeyErrorr   r   r(   rÑ   Ú
parse_attr)r?   r  rî   Ú
bound_argsr  rü   r  ÚsigkeysÚsigvalsr]   Úvr÷   Ú
constexprsÚpathÚattrvalsr  s                   rB   Ú
_pack_argszJITFunction._pack_argsŸ  s�  € à×'Ñ'¨Ó/ˆà#'§;¡;Ö/˜a�1—6“6Ð/ˆÐ/Ø!/Ö0˜A�1�Q“4Ð0ˆÐ0Ü(+¨G°WÓ(=×>™f˜q !�Q˜‘TÐ>ˆ	Ñ>à FÑ*ÐkÐ,kÓkÐ*Ø˜vÑ%ÐaÐ'aÓaÐ%Ø˜vÑ%ÐaÐ'aÓaÐ%Øò 	YˆAØ˜×(Ñ(Ò(¨Q°gÒ-=ÜÐSÐVWÑWÓXÐXð	Yô # 7Ñ,MÓNˆ
Ø[eÖfÐSW�dÔ-¬d°:×3DÑ3DÓ3FÓ.GÈÓNÑNÐfˆ
Ðfà"0Ö1˜Q�A�a“DÐ1ˆÐ1Ü˜hÑ(GÓHˆØPUÖVÈ1��G×&Ñ&Ô'8¸À1Ó'EÓFÑFÐVˆÐVà˜	 :¨uÐ4Ð4ùò% 0ùÚ0ùÛ>ùò gùâ1ùâVs"    D;¹E ÁEÃ +EÃ1EÄ#Ec               óà  — |j                  d| j                  «      xs t        j                  j                  |d<   t        j                  j
                  |d<   t        j                  j                  «       }t        j                  j                  |«      }| j                  D ]
  } ||i |¤Ž Œ | j                  |   \  }}	}
}} ||i |¤Ž\  }}}t        |	||«      }|j                  |d «      }|€4| j                  |||||«      \  }}}}| j                  |||||||«      }|€y t        «       }| j                   j#                  «       D ]6  \  \  }}\  }}|j                  ||«      x}|k7  sŒ$t%        d|› d|› d|› �«      ‚ |sê|€J ‚t'        |«      r ||«      }t)        |«      }|d   }|dkD  r|d   nd}|dkD  r|d   nd}t+        |d	«      r|j-                  «       } |j.                  ||g|j1                  «       ¢­Ž } |j2                  |||||j4                  |j6                  |t        j                  j8                  t        j                  j:                  g	|j1                  «       ¢­Ž  |S )
NÚdebugÚinstrumentation_moderT   z1 has changed since we compiled this kernel, from z to r   r   r   Úresult)rv   r¡  r   ÚruntimeÚcompilationr¢  r   rŒ  Úget_current_deviceÚget_current_streamr…  Údevice_cachesre  rŸ  Ú_do_compileÚobjectr=   rþ   rX   rm   r%   Úhasattrr£  Úlaunch_metadatarÑ   ræ   ÚfunctionÚpacked_metadataÚlaunch_enter_hookÚlaunch_exit_hook) r?   ré   rê   r‰   rî   rq  r’  r€  Úkernel_cacherd  r�   r  rŽ  r˜  r  rü   rý   Úkernelr÷   rœ  r  Únot_presentr2   r_   rp   Úglobals_dictÚnewValÚ	grid_sizeÚgrid_0Úgrid_1Úgrid_2r¬  s                                    rB   ræ   zJITFunction.run·  s›  € Ø Ÿ*™* W¨d¯j©jÓ9ÒP¼U¿]¹]×=PÑ=Pˆˆw‰Ü).×):Ñ):×)OÑ)OˆÐ%Ñ&ô —‘×1Ñ1Ó3ˆÜ—‘×1Ñ1°&Ó9ˆð ×&Ñ&ò 	"ˆDÙ�$Ð!˜&Ó!ð	"ð CG×BTÑBTÐU[ÑB\Ñ?ˆÐ&¨°¸ñ /5°dÐ.E¸fÑ.EÑ+ˆ
�N GäÐ 0°.À'ÓJˆØ×!Ñ! # tÓ,ˆð ˆ>Ø48·O±OÀGÈVÐU_ÐaoØDKó5MÑ1ˆG�Y 
¨Eð ×%Ñ% c¨9°f¸jÈ'ÐSXÐZ`ÓaˆFØˆ~Øô “hˆØ.2×.CÑ.C×.IÑ.IÓ.Kò 	qÑ*‰IˆT�1Ñ*˜˜\Ø&×*Ñ*¨4°Ó=Ð=�À#ÓEÜ"Ø& t fÐ,]Ð^aÐ]bÐbfÐgmÐfnÐoóqð qð	qñ
 àÐ#Ð#Ð#Ü˜Œ~Ù˜JÓ'�Ü˜D›	ˆIØ˜!‘WˆFØ )¨A¢�T˜!’W°1ˆFØ )¨A¢�T˜!’W°1ˆFÜ�v˜xÔ(ØŸ™›�à4˜f×4Ñ4°T¸6ÐXÀJ×DUÑDUÓDWÒXˆOØˆF�J‰J�v˜v v¨v°v·±È×H^ÑH^Ð`oÜ—}‘}×6Ñ6¼¿¹×8VÑ8VðnØYc×YjÑYjÓYlónàˆrC   c                óT   — | j                   €| j                  S | j                  |«      S rE   )Ú_reprr/  )r?   r_   s     rB   r{  zJITFunction.reprì  s"   € Ø $§
¡
Ð 2ˆt�}‰}ÐE¸¿
¹
À1»ÐErC   c	           	     óø  •— |r|ng }|r|ng }t         ‰| �  |«       |j                  | _        || _        || _        || _        || _        || _        g | _	        t        | j                  j                  j                  «       «      D ]T  \  }	}
|	|v xs |
j                  |v }|	|v xs |
j                  |v }| j                  j                  t!        |	|
||«      «       ŒV t#        | j$                  «      | _        d | _        || _        || _        | j                  D �cg c]  }|j                  ‘Œ c}| _        | j                  D �cg c]  }|j0                  sŒ|j2                  ‘Œ c}| _        g | _        y c c}w c c}w rE   )r0   r1   rJ   rQ   ÚversionrÈ   rÉ   r»  r¬  r  rg  r÷   r  rÑ   r2   r  rÄ   r   r�  r¨  r²  r¡  rU   Ú	arg_namesrÖ   rÆ   rœ  r…  )r?   r&  r½  rÈ   rÉ   r¡  rU   r{  r¬  rj  rÊ   ÚdnsÚdns_oaÚprA   s                 €rB   r1   zJITFunction.__init__ï  sW  ø€ á1BÑ-ÈÐÙKiÑ)GÐoqÐ&ä‰Ñ˜ÔØ—m‘mˆŒØˆŒØ!2ˆÔØ.LˆÔ+ØˆŒ
Ø.ˆÔàˆŒÜ! $§.¡.×";Ñ";×"BÑ"BÓ"DÓEò 	C‰HˆAˆuØÐ(Ð(ÒK¨E¯J©JÐ:KÐ,KˆCØÐ8Ð8Òh¸E¿J¹JÐJhÐ<hˆFØ�K‰K×Ñœ{¨1¨e°S¸&ÓAÕBð	Cô )¨×);Ñ);Ó<ˆÔð ˆŒØˆŒ
Ø ˆŒð +/¯+©+Ö6 Q˜!Ÿ&›&Ò6ˆŒØ*.¯+©+ÖH Q¸¿»˜1Ÿ5›5ÒHˆŒð  ˆÕùò	 7ùÚHs   ÄE2ÅE7ÅE7c           	     óÎ  — dd l }dd lm} t        j                  j                  «       }|j                  |«      }|d   | j                  k7  rt        d|d   › d| j                  › �«      ‚t        t        |d   «      }|d   }t        ||«      D ��	ci c]\  \  }}	||j                  j                  |	«      r|j                  |	«      n)t        |	t        «      rd|	v r|j!                  |	d   «      n|	“Œ^ }
}}	t        t        |d   «      }|d	   }t        t        ||«      «      }|d
   j#                  «       D ��	ci c]  \  }}	|t%        |	«      “Œ }}}	|d   j#                  «       D ��	ci c]#  \  }}	|t        |	t&        «      rt        |	«      n|	“Œ% }}}	|d   }| j(                  |   \  }}}}}|j+                  |«      }| j-                  ||||
||d¬«      S c c}	}w c c}	}w c c}	}w )Nr   r2   zSpecialization data is for z but trying to preload for rø   rù   rf   rú   rû   r÷   rü   rý   T)rê   )rÿ   Útriton.languager¸   r   rŒ  r¦  Úloadsr/  rX   rë   rb  r  rÀ   Úis_dtyper-   rB  rf   rþ   rh  r(   r¨  r•  r©  )r?   ry  rÿ   Útlrq  Údeserialized_objrø   rù   rý   rƒ   rœ  rú   rû   r  r÷   rü   r_   r  s                     rB   ÚpreloadzJITFunction.preload  s  € ÛÝ$Ü—‘×1Ñ1Ó3ˆØŸ:™:Ð&9Ó:ÐØ˜FÑ# t§}¡}Ò4ÜØ-Ð.>¸vÑ.FÐ-GÐGbÐcg×cpÑcpÐbqÐrótð täœEÐ#3°OÑ#DÓEˆØ(¨Ñ9ˆô
 " -°Ó?÷	
ñ ��Uð Ø!Ÿx™x×0Ñ0°Ô7ˆB�H‰H�UŒOÜ0:¸5Ä$Ô0GÈKÐ[`ÑL`ˆB�L‰L˜˜{Ñ+Ô,Ðfkñlð
ˆ
ñ 
ô œÐ 0°Ñ >Ó?ˆ
Ø% lÑ3ˆ
Ü”S˜ ZÓ0Ó1ˆð M]Ð]hÑLi×LoÑLoÓLq×r¹j¸cÀ5�SÔ2°5Ó9Ñ9Ðrˆ	Ñrð /¨yÑ9×?Ñ?ÓA÷
á��Uð ¤¨E´4Ô!8”�u”¸eÑCð
ˆñ 
ð ˜uÑ%ˆØ"×0Ñ0°Ñ8Ñˆˆ1ˆa�˜!Ø×'Ñ'¨Ó0ˆØ×ÑØØØØØØØð  ó 
ð 	
ùó'
ùó sùó
s   ÂA!GÄ8GÅ'(G!c           
     óB  ‡ ‡‡‡‡‡‡‡‡‡‡‡— ‰ j                   ‰   \  Š}Š}	}‰ j                  t        j                  j                  ‰‰‰‰‰‰g‰«      ry ‰ j                  ‰ ‰‰‰«      Št        j                  j                  «       }
|
�Ct        «       Št        ‰|	‰‰«      }ˆˆˆ ˆˆfd„}ˆˆˆˆˆˆˆ ˆˆf	d„}|
j                  |||«      }|S ‰ j                  ‰‰‰j                  ¬«      }|‰‰<   ‰ j                  t        j                  j                  ‰‰‰‰‰‰g‰«       |S )Nc                 óB   •— ‰j                  ‰‰‰j                  ‰ ¬«      S )N)r�   rü   Ú	_env_vars)r|  r   )Úenv_varsrü   r?   r@   r�   s   €€€€€rB   Úasync_compilez.JITFunction._do_compile.<locals>.async_compileG  s!   ø€ Ø—|‘| C°À×@PÑ@PÐ\d�|ÓeÐerC   c           
     ór   •	— | ‰‰<   ‰j                  t        j                  j                  ‰‰‰‰‰‰g‰	«       y rE   )rƒ  r   r¤  Újit_post_compile_hook)
r²  r  rœ  rq  r±  rý   rü   r?   r÷   rê   s
    €€€€€€€€€rB   Úfinalize_compilez1JITFunction._do_compile.<locals>.finalize_compileJ  s9   ø€ Ø$*�˜SÑ!Ø—‘¤§¡× CÑ CÀSÈ)ÐU[Ð]gÐipØ!& ¨õ1rC   )r�   rü   )r¨  rƒ  r   r¤  Újit_cache_hookr‰  r   Úactive_moderv   r   r   Úsubmitr|  r   rÏ  )r?   rý   r÷   rq  rœ  rü   r  rê   r_   r  Ú
async_moder[   rÍ  rÐ  r²  rÌ  r±  r@   r�   s   ````````       @@@@rB   r©  zJITFunction._do_compile:  s   ÿû€ Ø.2×.@Ñ.@ÀÑ.HÑ+ˆ�a˜ ¨!à�?‰?œ5Ÿ=™=×7Ñ7¸¸iÈÐQ[Ð]dÐglÐfmÐouÔvØØ�n‰n˜T 9¨j¸%Ó@ˆä#×/Ñ/×3Ñ3Ó5ˆ
ØÐ!ä6Ó8ˆHÜ% c¨7°G¸XÓFˆI÷fð f÷1ô 1ð
  ×&Ñ& y°-ÐAQÓRˆFð ˆð	 —\‘\ #¨f¸g×>NÑ>N�\ÓOˆFØ &ˆL˜ÑØ�O‰OœEŸM™M×?Ñ?ÀÀiÐQWÐYcÐelÐotÐnuØ"ô$àˆrC   c                ó   — t        d«      ‚)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rX   ©r?   r‰   rî   s      rB   Ú__call__zJITFunction.__call__W  s   € ÜÐWÓXÐXrC   c                óP   — d| j                   › d| j                  j                  › d�S )NzJITFunction(ú:r	  )rQ   r&  r¬   rG   s    rB   Ú__repr__zJITFunction.__repr__Z  s&   € Ø˜dŸk™k˜]¨!¨D¯G©G×,@Ñ,@Ð+AÀÐCÐCrC   )rª   zbool | None)NNNNNNN)rY   rJ   r¬   rn  rƒ  r†  r�  rŸ  ræ   r{  r1   rÈ  r©  r×  rÚ  r¯   r°   s   @rB   rZ  rZ  W  sZ   ø„ òð,
ð 
ó,
ò\(ò/ò5ò03òjFð mqØ;?õ" òH%
òNò:YöDrC   rZ  c                 ó   — y rE   rß   r%  s    rB   ÚjitrÜ  c  s   € àrC   ©r½  r{  r¬  rÈ   rÉ   r¡  rU   c                 ó   — y rE   rß   rÝ  s          rB   rÜ  rÜ  h  s   € ð rC   c               ó@   ‡‡‡‡‡‡‡— dˆˆˆˆˆˆˆfd„}| � || «      S |S )a<  
    Decorator for JIT-compiling a function using the Triton compiler.

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

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

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

    :param fn: the function to be jit-compiled
    :type fn: Callable
    c           
     ó¤   •— t        | «      sJ ‚t        j                  j                  rddlm}  || ‰‰‰‰‰‰‰¬«      S t        | ‰‰‰‰‰‰‰¬«      S )Nr   )ÚInterpretedFunction)r½  rÈ   rÉ   r¡  rU   r{  r¬  )rm   r   r¤  Ú	interpretÚinterpreterrá  rZ  )	r&  rá  r¡  rÈ   rÉ   r¬  rU   r{  r½  s	     €€€€€€€rB   Ú	decoratorzjit.<locals>.decorator“  sj   ø€ Ü˜Œ|Ðˆ|Ü�=‰=×"Ò"Ý8Ù& r°7ÐN_ØFdÐlqØ08¸tÐUdôfð fô ØØØ"3Ø/MØØ!ØØ /ô	ð 	rC   ©r&  r   rª   zJITFunction[T]rß   )	r&  r½  r{  r¬  rÈ   rÉ   r¡  rU   rä  s	    ``````` rB   rÜ  rÜ  v  s&   þ€ ÷:ó ð& 
€~Ù˜‹}Ðð ÐrC   c                  óN   — e Zd ZdZed„ «       Zdd„Zd„ Zed„ «       Zed„ «       Z	y)	rì   zr
    Can be used in place of real tensors when calling:
        kernel.warmup(MockTensor(torch.float32), ...)
    c                ól   — | j                   j                  dk(  r| j                  dk(  rt        | «      S | S )NrÀ   Útorch)rA   rY   rJ   rì   )rŠ   s    rB   rí   zMockTensor.wrap_dtype¸  s.   € à�=‰=×!Ñ! WÒ,°·±À7Ò1JÜ˜c“?Ð"Øˆ
rC   Nc                ó*   — |€dg}|| _         || _        y )Nr   )rÀ   Úshape)r?   rÀ   rê  s      rB   r1   zMockTensor.__init__¾  s   € Øˆ=Ø�CˆEØˆŒ
Øˆ�
rC   c                ó†   — dg}| j                   dd  D ]  }|j                  |d   |z  «       Œ t        t        |«      «      S )Nr   rµ   )rê  r  rb  Úreversed)r?   ÚstridesÚsizes      rB   ÚstridezMockTensor.strideÄ  sG   € Ø�#ˆØ—J‘J˜q˜r�Nò 	/ˆDØ�N‰N˜7 2™;¨Ñ-Õ.ð	/ä”X˜gÓ&Ó'Ð'rC   c                  ó   — y©Nr   rß   rß   rC   rB   Údata_ptrzMockTensor.data_ptrÊ  ó   € àrC   c                  ó   — yrñ  rß   rß   rC   rB   Ú	ptr_rangezMockTensor.ptr_rangeÎ  ró  rC   rE   )
rY   rJ   r¬   r­   Ústaticmethodrí   r1   rï  rò  rõ  rß   rC   rB   rì   rì   ²  sM   „ ñð
 ñó ðó
ò(ð ñó ðð ñó ñrC   rì   c                  óJ   — e Zd Zd„ Zd„ Zd„ Zdd„Zd„ Zd„ Zd„ Z	d„ Z
d	„ Zd
„ Zy)ÚTensorWrapperc                óš   — || _         || _        |j                  | _        |j                  | _        | j                  j                  | _        y rE   )rÀ   ÚbaseÚdatarq  rê  )r?   rú  rÀ   s      rB   r1   zTensorWrapper.__init__Õ  s5   € ØˆŒ
ØˆŒ	Ø—I‘IˆŒ	Ø—k‘kˆŒØ—Y‘Y—_‘_ˆ�
rC   c                ó6   — | j                   j                  «       S rE   )rú  rò  rG   s    rB   rò  zTensorWrapper.data_ptrÜ  s   € Ø�y‰y×!Ñ!Ó#Ð#rC   c                ó4   —  | j                   j                  |Ž S rE   )rú  rï  )r?   r‰   s     rB   rï  zTensorWrapper.strideß  s   € Øˆt�y‰y×Ñ Ð&Ð&rC   c                ó<   — d| j                   › d| j                  › d�S )NzTensorWrapper[rp  r	  )rÀ   rú  rG   s    rB   Ú__str__zTensorWrapper.__str__â  s   € Ø §
¡
˜|¨2¨d¯i©i¨[¸Ð:Ð:rC   c                ó6   — | j                   j                  «       S rE   )rú  Úelement_sizerG   s    rB   r  zTensorWrapper.element_sizeå  s   € Ø�y‰y×%Ñ%Ó'Ð'rC   c                ó^   — t        | j                  j                  «       | j                  «      S rE   )rø  rú  ÚcpurÀ   rG   s    rB   r  zTensorWrapper.cpuè  s   € Ü˜TŸY™YŸ]™]›_¨d¯j©jÓ9Ð9rC   c                óN   — | j                   j                  |j                   «       y rE   )rú  Úcopy_)r?   Úothers     rB   r  zTensorWrapper.copy_ë  s   € Ø�	‰	�‰˜Ÿ
™
Õ#rC   c                ó^   — t        | j                  j                  «       | j                  «      S rE   )rø  rú  ÚclonerÀ   rG   s    rB   r  zTensorWrapper.cloneî  s   € Ü˜TŸY™YŸ_™_Ó.°·
±
Ó;Ð;rC   c                ó`   — t        | j                  j                  |«      | j                  «      S rE   )rø  rú  ÚtorÀ   )r?   rq  s     rB   r
  zTensorWrapper.toñ  s    € Ü˜TŸY™YŸ\™\¨&Ó1°4·:±:Ó>Ð>rC   c                ó`   — t        | j                  j                  |«      | j                  «      S rE   )rø  rú  Ú	new_emptyrÀ   )r?   Úsizess     rB   r  zTensorWrapper.new_emptyô  s"   € Ü˜TŸY™Y×0Ñ0°Ó7¸¿¹ÓDÐDrC   NrÞ   )rY   rJ   r¬   r1   rò  rï  rÿ  r  r  r  r  r
  r  rß   rC   rB   rø  rø  Ó  s5   „ ò%ò$ò'ó;ò(ò:ò$ò<ò?óErC   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 r$  )r-   rø  rú  rÀ   r«  r¡   rj   )ÚtensorrÀ   s     rB   Úreinterpretr  ø  sk   € Ü�&œ-Ô(Ø�F—K‘K×%Ñ%Ò%à—;‘;Ðô ! §¡¨eÓ4Ð4Ü	�˜Ô	$ä˜V UÓ+Ð+äÐ/´°V³¨~¸QÐ?Ó@Ð@rC   c                óL  — | }t        |t        «      s|j                  }t        |t        «      sŒ|j                  j                  j                  }|j
                  }t        |j                  «      D ].  \  }}|j                  «       j                  d«      sŒ&||z  } ||fS  ||fS )Nzdef )
r-   rV   r&  Ú__code__Úco_filenamer,  rg  r+  rº   rO   )r&  Úbase_fnÚ	file_nameÚ
begin_lineÚidxÚlines         rB   Úget_jit_fn_file_liner    s�   € Ø€GÜ˜¤+Ô.Ø—*‘*ˆô ˜¤+Õ.à—
‘
×#Ñ#×/Ñ/€IØ×-Ñ-€Jô ˜wŸ™Ó/ò ‰	ˆˆTØ�:‰:‹<×"Ñ" 6Õ*Ø˜#ÑˆJØØ�jÐ Ð ð	ð �jÐ Ð rC   c                  ó(   — e Zd Zd„ Zed„ «       Zd„ Zy)ÚBoundConstexprFunctionc                ó    — || _         || _        y rE   )Ú__self__Ú__func__)r?   Úinstancer&  s      rB   r1   zBoundConstexprFunction.__init__  s   € Ø ˆŒØˆ�rC   c                ó.   — | j                   j                  S rE   )r  r[   rG   s    rB   r[   z BoundConstexprFunction.cache_key  s   € à�}‰}×&Ñ&Ð&rC   c                óB   —  | j                   | j                  g|¢­i |¤ŽS rE   )r  r  rÖ  s      rB   r×  zBoundConstexprFunction.__call__#  s!   € Øˆt�}‰}˜TŸ]™]Ð<¨TÒ<°VÑ<Ð<rC   N)rY   rJ   r¬   r1   r®   r[   r×  rß   rC   rB   r  r    s    „ òð ñ'ó ð'ó=rC   r  c                  ó0   ‡ — e Zd Zˆ fd„Zd„ Zddœd„Zˆ xZS )ÚConstexprFunctionc                ó$   •— t         ‰| �  |«       y rE   )r0   r1   )r?   r&  rA   s     €rB   r1   zConstexprFunction.__init__)  s   ø€ Ü‰Ñ˜ÕrC   c                ó"   — |�t        || «      S | S rE   )r  )r?   r  Úobjclasss      rB   Ú__get__zConstexprFunction.__get__,  s   € àˆ?Ü)¨#¨tÓ4Ð4ØˆrC   N)Ú	_semanticc               ó  — ddl m}m} |D �cg c]
  } ||«      ‘Œ }}|j                  «       D ��ci c]  \  }}| ||«      “Œ }}} | j                  |i |¤Ž}	|€|	S t
        j                  j                  r|	S  ||	«      S c c}w c c}}w )Nr   )Ú_unwrap_if_constexprrf   )r·   r*  rf   rþ   r&  r   r¤  râ  )
r?   r(  r‰   rî   r*  rf   r  r]   r›  Úress
             rB   r×  zConstexprFunction.__call__2  sŽ   € ßHà15Ö6¨AÑ$ QÕ'Ð6ˆÐ6Ø;A¿<¹<»>×J±°!°Q�!Ñ)¨!Ó,Ñ,ÐJˆÑJð ˆd�g‰g�tÐ&˜vÑ&ˆàÐàˆJô �=‰=×"Ò"ØˆJÙ˜‹~Ðùò 7ùÛJs
   �B±B)rY   rJ   r¬   r1   r'  r×  r¯   r°   s   @rB   r#  r#  '  s   ø„ ôòð )-÷ ð rC   r#  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#  r%  s    rB   Úconstexpr_functionr-  E  s   € ô ˜RÓ Ð rC   rÞ   )Frå  )r{  úOptional[Callable]r¬  r.  rÈ   úOptional[Iterable[int | str]]rÉ   r/  r¡  úOptional[bool]rU   r0  rª   zCallable[[T], JITFunction[T]]rE   )r&  zOptional[T]r{  r.  r¬  r.  rÈ   r/  rÉ   r/  r¡  r0  rU   r0  rª   zKernelInterface[T])FÚ
__future__r   r   ry   r.   r3   rL   r’   r0  r5  r3  Úcollectionsr   Údataclassesr   Ú	functoolsr   Útypingr   r	   r
   r   r   r   r   r   r   Útriton.backendsr   Útypesr   rK   r   r   r   Ú_utilsr   r   r   r   Úcacher   Útriton._C.libtritonr   r   r;   r:   r   ÚNodeVisitorr!   r¼   rÄ   rã   rå   r  r"  r'  rV   rY  re  rh  rZ  rÜ  rì   rø  r  r  r  r#  r-  rß   rC   rB   ú<module>r<     s0  ðß ,Û 
Û Û Û Û Û Û 	Û Ý #Ý !Ý %ß ]× ]Õ ]å 'Ý Ý Ý Ý ß `Ó `Ý  ß Wà!€Ø3€áˆCƒL€ôg!˜Ÿ™ô g!ó^D÷4/>ñ />ódTôZ�g˜a‘jô Zò&ò"9*òx0÷b1ñ b1ðJ ÷ð ó ðòò0	ôDD�+˜¨qÑ1ô DDðX 
òó 
ðð 
ð Ø#Ø*.Ø7;ØDHØ Ø#ñ
ð ð
ð (ð	
ð
 5ð
ð %Bð
ð ð
ð ð
ð #ò
ó 
ð
ð ð4ð Ø#Ø*.Ø7;ØDHØ Ø#ñ4Øð4ð ð	4ð
 (ð4ð 5ð4ð %Bð4ð ð4ð ð4ð ó4÷xñ ÷B"Eñ "EòJAò!ô$=˜[ô =ô˜ô ó<!rC   