Ë
    1²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 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,m-Z- dZ.dZ/ ed«      Z0 G d„ dejb                  «      Z2d6d„Z3 G d„ d«      Z4d7d„Z5 G d„ dee0   «      Z6d„ Z7d„ Z8d„ Z9 G d„ d «      Z:i Z;e G d!„ d"«      «       Z<d#„ Z=d$„ Z> G d%„ d&e:e6e0   «      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,«      ZA G d-„ d.«      ZBd/„ ZCd0„ ZD G d1„ d2e:«      ZE G d3„ d4e:«      ZFd5„ ZGy)<é    )Ú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_implÚirz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)Úselfr3   r8   r9   ÚsrcÚ	__class__s        €úQ/var/www/html/ia-enhanced/venv/lib/python3.12/site-packages/triton/runtime/jit.pyr2   zDependenciesFinder.__init__.   st   ø€ Ü‰ÑÔØˆŒ	Ü—n‘n S§Z¡Z°Ó%8Ó9ˆŒð ˆŒØ"ˆŒò*
ˆÔ&ô ÜØØð	"
ˆÔð" TVˆÔà*/ˆÕ'ó    c                ó6   — | j                   j                  «       S ©N)r7   Ú	hexdigest©r@   s    rC   ÚretzDependenciesFinder.retY   s   € à�{‰{×$Ñ$Ó&Ð&rD   c                ó†   — t        j                  |j                  «      ryt        |dd«      }|j	                  t
        «      S )NTÚ
__module__Ú )ÚinspectÚ	isbuiltinÚfuncr-   Ú
startswithr<   )r@   ÚnoderO   Úmodules       rC   Ú_is_triton_builtinz%DependenciesFinder._is_triton_builtin]   s6   € Ü×Ñ˜TŸY™YÔ'ØÜ˜˜|¨RÓ0ˆØ× Ñ ¤Ó/Ð/rD   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ÚRuntimeErrorr3   Ú__name__ÚupdateÚ	cache_keyÚstrr-   r7   r6   )r@   rO   ÚkÚvar_nameÚ_Úv1Úv2Úfunc_keys           rC   Ú_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Õ4rD   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__rK   rL   ztriton.language.extra.libdevicez!Unsupported function referenced: )Úlanguage.corerg   Útyper   r-   Ú
hash_attrsÚrecord_referencer.   rW   rd   ÚcallablerY   r?   r/   Údeepcopyr>   Úid)r@   ÚvalÚvar_dictr3   rg   Úattrs         rC   rm   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ØrD   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)r8   Úgetr9   )r3   rq   r@   s     €rC   Úname_lookupz2DependenciesFinder.visit_Name.<locals>.name_lookup¤   sZ   ø€ Ø—,‘,×"Ñ" 4¨Ó.ˆCØˆØ˜DŸL™LÐ(Ð(Ø—.‘.×$Ñ$ T¨4Ó0ˆCØˆØ˜DŸN™NÐ*Ð*ØrD   )rk   ÚctxÚastÚStorerp   Úlocal_namesr:   rm   )r@   rQ   rx   rq   rr   s   `    rC   Ú
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Øˆ
rD   c                ó^   — |j                   D �cg c]  }| j                  |«      ‘Œ c}S c c}w rF   )ÚeltsÚvisit)r@   rQ   Úelts      rC   Ú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 )NrZ   rL   )	r€   Úvaluer.   rz   Ú	Attributer-   r=   rs   rm   )r@   rQ   ÚlhsÚlhs_namerI   s        rC   Ú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Ô"Øˆ
rD   c                ó–   — |j                   j                   D �ch c]  }|j                  ’Œ c}| _        | j                  |«       y c c}w rF   )ÚargsÚargr|   Úgeneric_visit)r@   rQ   r‹   s      rC   Ú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     €rC   Ú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@   rQ   r’   r‹   s   `   rC   Ú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—}‘}Õ%rD   c                óº   — | j                  |«      }t        |t        «      r| xj                  t	        |«      z  c_        y | j                  j                  |«       y rF   )r€   r.   r)   r|   ÚsetÚadd)r@   rQ   Útargets      rC   ÚvisitAssnTargetz"DependenciesFinder.visitAssnTargeté   sE   € ð —‘˜DÓ!ˆÜ�fœdÔ#Ø×Ò¤ F£Ñ+Öà×Ñ× Ñ  Õ(rD   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@   rQ   s     rC   Úvisit_AssignzDependenciesFinder.visit_Assignò   sG   € Üˆt�|‰|Ó Ò!ô
 ÐPÓQÐQà×Ñ˜TŸ\™\¨!™_Ô-ð 	×Ñ˜4Õ rD   c                ó\   — | j                  |j                  «       | j                  |«       y rF   ©rŸ   rž   rŒ   r£   s     rC   Úvisit_AnnAssignz"DependenciesFinder.visit_AnnAssignÿ   ó$   € Ø×Ñ˜TŸ[™[Ô)ð 	×Ñ˜4Õ rD   c                ó\   — | j                  |j                  «       | j                  |«       y rF   r¦   r£   s     rC   Ú	visit_ForzDependenciesFinder.visit_For  r¨   rD   )ÚreturnÚNonerv   )rZ   rK   Ú__qualname__Ú__doc__r2   ÚpropertyrI   rS   rd   rm   r}   r‚   rˆ   r�   rš   rŸ   r¤   r§   rª   Ú__classcell__©rB   s   @rC   r"   r"   "   s`   ø„ ñ	õ)0ðV ñ'ó ð'ò0ò5ó$%òNò06ò
	ò!ò
&ò@)ò!ò!ö!rD   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.Ú_trL   )Útriton.language.coreÚlanguageÚcorer.   r]   ÚstriprP   ÚremoveprefixÚ_normalize_tyÚendswithÚpointer_typeÚ
element_tyÚdtyper3   rk   rZ   r   rw   Úreplace)Útyrº   s     rC   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ÐCrD   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 rF   )ÚnumÚ_paramÚdo_not_specializeÚdo_not_specialize_on_alignment)r@   rÇ   ÚparamrÉ   rÊ   s        rC   r2   zKernelParam.__init__.  s    € àˆŒØˆŒØ!2ˆÔØ.LˆÕ+rD   c                ó.   — | j                   j                  S rF   )rÈ   r3   rH   s    rC   r3   zKernelParam.name5  s   € à�{‰{×ÑÐrD   c                óÐ   — | j                   j                  r1| j                   j                  t        j                  j                  k(  ryt        | j                   j                  «      S )NrL   )rÈ   Ú
annotationrM   Ú	ParameterÚemptyr½   rH   s    rC   rÎ   zKernelParam.annotation9  sD   € à�{‰{×%Ò%¨¯©×)?Ñ)?Ä7×CTÑCT×CZÑCZÒ)ZØÜ˜TŸ[™[×3Ñ3Ó4Ð4rD   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   rL   )rÎ   rP   rœ   r   Úvalues)r@   Úas     rC   Úannotation_typezKernelParam.annotation_type?  s]   € à�O‰OˆØ�<‰<˜ÔØ�!�"�‰AØ�\‰\˜#ÔØ�!�"�ˆAØ”Ô.×5Ñ5Ó7Ó8Ñ8Ø—?‘?Ð"ØrD   c                ó   — d| j                   v S ©Nrg   )rÎ   rH   s    rC   Úis_constexprzKernelParam.is_constexprJ  s   € à˜dŸo™oÐ-Ð-rD   c                ór   — | j                   ryd| j                  v xs | j                  j                  d«      S )NFr³   rµ   )r×   rÎ   rP   rH   s    rC   Úis_constzKernelParam.is_constN  s1   € à×ÒØØ˜$Ÿ/™/Ð)ÒM¨T¯_©_×-GÑ-GÈÓ-MÐMrD   c                ó.   — | j                   j                  S rF   )rÈ   ÚdefaultrH   s    rC   rÛ   zKernelParam.defaultT  s   € à�{‰{×"Ñ"Ð"rD   c                ód   — | j                   j                  t        j                  j                  k7  S rF   )rÈ   rÛ   rM   rÏ   rÐ   rH   s    rC   Úhas_defaultzKernelParam.has_defaultX  s#   € à�{‰{×"Ñ"¤g×&7Ñ&7×&=Ñ&=Ñ=Ð=rD   N)rÇ   r%   rË   zinspect.ParameterrÉ   ÚboolrÊ   rÞ   ©r«   r]   )rZ   rK   r­   r®   r2   r   r3   rÎ   rÔ   r×   rÙ   r¯   rÛ   rÝ   © rD   rC   rÅ   rÅ   +  s¢   „ ÙLðMØ15óMð ñ ó ð ð ò5ó ð5ð
 òó ðð ñ.ó ð.ð ñNó ðNð
 ñ#ó ð#ð ñ>ó ñ>rD   rÅ   c                ó6   — d}d}t        t        | |||«      d   S )NFTr   )r   r   )r‹   Ú
specializerÙ   Úaligns       rC   Úmangle_typerä   ]  s%   € Ø€HØ€EÜ!¤+¨s°H¸jÈ%ÓPÐQRÑSÐSrD   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       rC   rë   zKernelInterface.warmupf  s*   € Øˆt�x‰x´´Z×5JÑ5JÈDÓ1Q˜T¨$Ñ\ÐU[Ñ\Ð\rD   c               ó   — t        d«      ‚)Nzrun not implemented)ÚNotImplementedError)r@   rê   rë   rŠ   rï   s        rC   rç   zKernelInterface.runi  s   € Ü!Ð"7Ó8Ð8rD   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     €€rC   ú<lambda>z-KernelInterface.__getitem__.<locals>.<lambda>r  s   ø€  x t§x¡xÈ$°TÀ%Ñ'YÐRXÑ'Y€ rD   rà   )r@   rê   s   ``rC   Ú__getitem__zKernelInterface.__getitem__l  s   ù€ ô ZÐYrD   N)r«   r    )rZ   rK   r­   Ú__annotations__rë   rç   rõ   rà   rD   rC   ræ   ræ   c  s   … Ø	
ƒFò]ò9ôZrD   ræ   c           
     óÜ  — |j                  «       D ��ci c]“  \  }}||j                  j                  dk(  rt        |«      nh|j                  j                  dk(  rd|j                  inA|j                  j                  dk(  r'd|j
                  › d|j                  j                  › �in|“Œ• }}}dd l}| ||j                  «       D �	cg c]  }	t        |	«      ‘Œ c}	t        |j                  «       «      |j                  «       D �	cg c]  }	t        |	«      ‘Œ c}	t        |j                  «       «      |j                  |j                  dœ	}
|j                  |
«      }|S c c}}w c c}	w c c}	w )NrÁ   rg   ÚJITFunctionÚjit_functionú:r   )	r3   Ú	signatureÚconstant_keysÚconstant_valsÚ
attrs_keysÚ
attrs_valsÚoptionsÚkeyrž   )ÚitemsrB   rZ   r]   r„   rR   Úfnr­   ÚjsonrX   r)   rÒ   Ú__dict__Údumps)r3   rû   Ú	constantsÚattrsr   r  rž   r„   r  ÚxÚobjÚserialized_objs               rC   Úserialize_specialization_datar  v  sI  € ð
 $Ÿ/™/Ó+÷	ñ ˆC�ð 	˜5Ÿ?™?×3Ñ3°wÒ>ŒS�ŒZØ�?‰?×#Ñ# {Ò2ð FQÐRW×R]ÑR]ÑD^à�?‰?×#Ñ# }Ò4ð :HÈEÏLÉLÈ>ÐYZÐ[`×[cÑ[c×[pÑ[pÐZqÐIrÑ8sØ:?ñ	@ð€Iñ ó à 9ÐQZ×Q_ÑQ_ÓQaÖ?bÈAÄÀQÅÒ?bÜˆY×ÑÓÓ À%Ç*Á*Ã,Ö0O¸Q´°aµÒ0OÔ_cÐdi×dpÑdpÓdrÓ_sØ×#Ñ#¨C¸6¿?¹?ñ€Cð
 —Z‘Z “_€NØÐùóùò @cùÚ0Os   ”BEÃE$ÄE)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Û   rM   rÏ   rÐ   )r	  s    rC   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€ rD   z
def dynamic_func(z	**optionsz):
    params = {ú'z': z}
    specialization = [ú,z-]
    return params, specialization, options
Údefault_Úspecialize_implÚbackendrW   Údynamic_func)r&   Ú
parametersÚziprX   r×   ÚappendrÙ   rÉ   rÊ   rÔ   r.   r]   Újoinr)   rì   r  rÛ   rM   rÏ   rÐ   r   rW   Úexec)ÚsigÚkparamsr  Úspecializationr3   ÚkprÙ   râ   rã   rI   r‹   Ú	func_bodyrË   Úfunc_namespacer  s                  rC   Ú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ú.)rK   r­   ©r  s    rC   Úget_full_namer+  Ä  s   € Ø�m‰mˆ_˜A˜bŸo™oÐ.Ð/Ð/rD   c                  ót   — 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y)rW   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 filerL   z^def\s+\w+\s*\()r  rM   rû   ÚgetsourcelinesÚraw_srcÚstarting_line_numberÚOSErrorÚ
ValueErrorr+  Ú_fn_nameÚ	threadingÚRLockÚ
_hash_lockÚtextwrapÚdedentr  ÚreÚsearchÚ	MULTILINEÚstartÚ_srcÚhashr>   r®   rZ   r­   Ú__globals__rK   )r@   r  ÚerA   s       rC   r2   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                ó  — | j                   }|j                  €| j                  S t        |j                  j
                  |j                  «      D ��ci c]  \  }}||j                  “Œ }}}| j                  |z  S c c}}w rF   )r  Ú__closure__r?  r  Ú__code__Úco_freevarsÚcell_contents)r@   r  r3   Úcellr9   s        rC   Úget_capture_scopezJITCallable.get_capture_scopeì  sr   € Ø�W‰WˆØ�>‰>Ð!Ø×#Ñ#Ð#Ü@CÀBÇKÁK×D[ÑD[Ð]_×]kÑ]kÓ@l×m±*°$¸�T˜4×-Ñ-Ñ-Ðmˆ	ÑmØ×Ñ )Ñ+Ð+ùó ns   ÁA;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:)r3   r8   r9   rA   r   rf   r$   )r6  r>  r3  rM   Úgetclosurevarsr  r9   r"   r?  rA   r€   ÚparserI   r]   r0  ÚdictÚsortedr>   r  r¸   rg   r.   r4   r5   r6   rG   )r@   r9   Údependencies_finderrg   r3   r`   rq   s          rC   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 rF   )r>  r\   rH   s    rC   Ú__hash__zJITCallable.__hash__
  s   € Ü�D—N‘NÓ#Ð#rD   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   )rz   rJ  r=  r.   ÚModuler&   ÚbodyÚFunctionDef)r@   Útrees     rC   rJ  zJITCallable.parse  s_   € Ü�y‰y˜Ÿ™Ó#ˆÜ˜$¤§
¡
Ô+Ð+Ð+Ü�4—9‘9‹~ Ò"Ð"Ð"Ü˜$Ÿ)™) A™,¬¯©Ô8Ð8Ð8ØˆrD   c                ó   — ddl m}  || «      S )Nr   )Úconstexpr_type)r¸   rV  )r@   rV  s     rC   rk   zJITCallable.type  s   € å7Ù˜dÓ#Ð#rD   c                 ó   — y rF   rà   )r@   Úhandless     rC   Ú_flatten_irzJITCallable._flatten_ir  s   € ØrD   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>  r=  )r@   Únew_srcs     rC   Ú_unsafe_update_srczJITCallable._unsafe_update_src  s   € ð ˆŒ	Øˆ�	rD   c                ó   — t        d«      ‚)NzqCannot set attribute 'src' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)ÚAttributeErrorrH   s    rC   Ú_set_srczJITCallable._set_src)  s   € Üð (ó )ð 	)rD   c                ó   — | j                   S rF   )r=  rH   s    rC   Ú_get_srczJITCallable._get_src.  s   € Ø�y‰yÐrD   )ÚfgetÚfsetNrß   )rX  zlist[ir.value]r«   r¬   )rZ   rK   r­   r2   rG  r¯   r\   rO  rJ  rk   rY  r\  r_  ra  rA   rà   rD   rC   rW   rW   È  sa   „ ò (òD,ð òó ðò,$òð ñ$ó ð$óòò)ò
ñ ˜ xÔ
0�CrD   rW   c                  ó,   — e Zd ZU ded<   ded<   ded<   y)ÚJitFunctionInfor   rR   r]   r3   rø   rù   N)rZ   rK   r­   rö   rà   rD   rC   re  re  7  s   … àÓØ
ƒIØÔrD   re  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rF   rà   )Ú.0r‹   Úreplace_callabless     €rC   ú	<genexpr>z?compute_cache_key.<locals>.replace_callables.<locals>.<genexpr>L  s   øè ø€ Ò?°CÑ*¨3×/Ñ?ùs   ƒ)r.   r)   r   rB   ÚtuplerW   r\   )r
  r‹   Úresultsrj  s      €rC   rj  z,compute_cache_key.<locals>.replace_callablesE  s’   ø€ Ü�cœ4Ô Ø69Ö:¨sÑ% cÕ*Ò:Ð:Ü˜3ÔØ9<Ö=°#Ñ(¨Õ-Ð=ˆGÐ=Ø �3—=‘= 'Ð*Ð*Ü˜œUÔ#ÜÓ?¸3Ô?Ó?Ð?Ü˜œ[Ô)Ø—=‘=Ð Øˆ
ùò ;ùâ=s
   –B·B)rl  r]   rw   )Úkernel_key_cacher#  r   r  r\   rj  s        @rC   Úcompute_cache_keyro  >  se   ø€ Ü�Ó ¤# g£,Ð
/€CØ ×$Ñ$ S¨$Ó/€IØÐØÐô
ô Ñ% nÓ5Ó6¼¸W»ÑE€IØ%Ð�SÑØÐrD   c                ó~   — t        | t        «      s| S t        | «      D ]  \  }}t        |«      | |<   Œ t	        | «      S rF   )r.   r)   Ú	enumerateÚconvert_to_tuple_if_listrl  )ÚitemÚiÚnested_values      rC   rr  rr  V  sD   € ä�dœDÔ!Øˆô % T›?ò 9‰ˆˆ<Ü*¨<Ó8ˆˆQŠð9ô �‹;ÐrD   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 )rø   c                 ó   — y)NFrà   rH   s    rC   Úis_gluonzJITFunction.is_gluond  s   € ØrD   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­   rK   r  r  Úparamsr3   r|  r}  r~  r  r€  r+  r  r�  re  )r@   Úhookr  rû   rž   r{  r  r   r‚  r„  r3   rR   rË   rÃ   Ú	arg_reprsr…  Ú	full_namerƒ  rï   s                      rC   Ú
_call_hookzJITFunction._call_hookg  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Ø<BóDÐð #ØØ"Ø ×*Ñ*Ø×(Ñ(Ø!×,Ñ,Ø '× 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)rn   Úpre_run_hooksr  )r@   rŠ  s     rC   Úadd_pre_run_hookzJITFunction.add_pre_run_hook—  s$   € ô
 ˜Œ~Ðˆ~Ø×Ñ×!Ñ! $Õ'rD   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           rC   Úcreate_binderzJITFunction.create_binderŸ  sd   € ÷ 	PÓOÜ—‘×1Ñ1Ó3ˆÙ˜vÓ&ˆØ,ˆÔØˆŒØ"ˆŒÜ/°·±ÀÇÁÈWÓUˆØ�2�v˜w¨Ð.Ð.rD   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   d
k(  rdn|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 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`   rq   s     rC   rô   z(JITFunction._pack_args.<locals>.<lambda>»  s   € ¸3À+Ñ;M€ rD   rg   rL   r   c                ó"   — t        |t        «      S rF   )r.   r]   )r`   r	  s     rC   rô   z(JITFunction._pack_args.<locals>.<lambda>¿  s   € ´ZÀÄ3Ó5G€ rD   )Úparse_optionsr‰  r3   r  r  ÚKeyErrorr   r   r)   rÒ   Ú
parse_attr)r@   r  rï   Ú
bound_argsr#  r   r	  ÚsigkeysÚsigvalsr^   Úvrû   Ú
constexprsÚpathÚattrvalsr  s                   rC   Ú
_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àAOÖP¸A˜!˜A™$ +Ò-‘B°1°Q±4Ñ7ÐPˆÐPÜ˜hÑ(GÓHˆØPUÖVÈ1��G×&Ñ&Ô'8¸À1Ó'EÓFÑFÐVˆÐVà˜	 :¨uÐ4Ð4ùò% 0ùÚ0ùÛ>ùò gùâPùâVs"    E¹E
ÁEÃ +EÃ1EÄ#Ec               óH  — |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                  j                  �6t        j                  j                  «       \  }}|j                  d|› d�«       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	}  |j.                  ||g|j1                  «       ¢­Ž }! |j2                  ||| ||j4                  |j6                  |!t        j                  j8                  t        j                  j:                  g	|j1                  «       ¢­Ž  |S )NÚdebugÚinstrumentation_modez("custom_pipeline", r  rU   z1 has changed since we compiled this kernel, from z to r   r   r   )rw   r«  r   ÚruntimeÚcompilationr¬  r   r–  Úget_current_deviceÚget_current_streamr�  Údevice_cachesÚadd_stages_inspection_hookr  ro  r©  Ú_do_compileÚobjectr>   r  rY   rn   r&   Úlaunch_metadatarÒ   rç   ÚfunctionÚpacked_metadataÚlaunch_enter_hookÚlaunch_exit_hook)"r@   rê   rë   rŠ   rï   r{  rœ  rŠ  Úkernel_cachern  rž   r  r˜  r¢  r#  r   Úinspect_stages_keyÚinspect_stages_hashr  Úkernelrû   r¦  r  Únot_presentr3   r`   rq   Úglobals_dictÚnewValÚ	grid_sizeÚgrid_0Úgrid_1Úgrid_2rµ  s"                                     rC   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ô �=‰=×3Ñ3Ð?Ü6;·m±m×6^Ñ6^Ó6`Ñ3ÐÐ 3Ø×!Ñ!Ð$8Ð9LÐ8MÈQÐ"OÔPäÐ 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à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àˆrD   c                óT   — | j                   €| j                  S | j                  |«      S rF   )Ú_reprr3  )r@   r`   s     rC   r…  zJITFunction.reprý  s"   € Ø $§
¡
Ð 2ˆt�}‰}ÐE¸¿
¹
À1»ÐErD   c	           	     óP  •— |r|ng }|r|ng }t         ‰| �  |«       |j                  | _        || _        || _        || _        || _        || _        | t        | j                  › d| j                  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]  }|j6                  sŒ|j8                  ‘Œ c}| _        g | _        y c c}w c c}w )Nrú   )r1   r2   rK   rR   ÚversionrÉ   rÊ   rÆ  rµ  Ú_triton_jit_function_registryr  r­   r‰  rq  rû   r  rÒ   r3   r  rÅ   r   r™  r±  r½  r«  rV   Ú	arg_namesr×   rÇ   r¦  r�  )r@   r  rÈ  rÉ   rÊ   r«  rV   r…  rµ  rt  rË   ÚdnsÚdns_oaÚprB   s                 €rC   r2   zJITFunction.__init__   s|  ø€ á1BÑ-ÈÐÙKiÑ)GÐoqÐ&ä‰Ñ˜ÔØ—m‘mˆŒØˆŒØ!2ˆÔØ.LˆÔ+ØˆŒ
Ø.ˆÔàQUÔ%¨¯©¨°Q°t·w±w×7KÑ7KÐ6LÐ&MÑNàˆŒÜ! $§.¡.×";Ñ";×"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   ÅFÅ.F#Æ F#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   }| j                  |   \  }}}}	}|d   }
|j                  |
k7  rt        d|
› d|› �«      ‚ˆfd„}t        ||«      D ��ci c]  \  }}| ||«      “Œ }}}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)                  ||||||d¬«      S c c}}w c c}}w c c}}w )Nr   r3   zSpecialization data is for z but trying to preload for rü   rý   rž   c                ó  •— ‰j                   j                  | «      r‰j                  | «      S t        | t        «      rAd| v r‰j	                  | d   «      S d| v r%| d   }|t
        v r	t
        |   S t        d|› d�«      ‚| S )Nrg   rù   zUnable to resolve JITFunction z for preload)rÁ   Úis_dtyper.   rK  rg   rÉ  rY   )r„   Újf_keyÚtls     €rC   Ú_decode_constantz-JITFunction.preload.<locals>._decode_constant7  s‰   ø€ Ø�x‰x× Ñ  Ô'Ø—x‘x “Ð&Ü˜%¤Ô&Ø %Ñ'ØŸ<™<¨¨kÑ(:Ó;Ð;Ø! UÑ*Ø" >Ñ2�FØÔ!>Ñ>Ü<¸VÑDÐDÜ&Ð)GÈÀxÈ|Ð'\Ó]Ð]ØˆLrD   rþ   rÿ   rû   r   r  T)rë   )r  Útriton.languager¹   r   r–  r¯  Úloadsr3  rY   rì   rl  r±  r  r  rK  r  rr  r.   r)   rŸ  r³  )r@   rƒ  r  r{  Údeserialized_objrü   rý   r`   rž   r  Údeserialized_targetrÓ  r  r„   r¦  rþ   rÿ   r  rû   r   rÒ  s                       @rC   Ú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ˆØ#'×#5Ñ#5°fÑ#=Ñ ˆˆ1ˆf�g˜qØ.¨xÑ8Ðð �?‰?Ð1Ò1ÜÐ!<Ð=PÐ<QÐQlÐmsÐltÐuÓvÐvô	ô FIÈÐXeÓEf×g±z°s¸E�cÑ+¨EÓ2Ñ2Ðgˆ
ÑgÜœÐ 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ˆØ×ÑØØØØØØØð  ó 
ð 	
ùó hùó sùó
s   ÃF2Ä*F8Å(F>c                óR  ‡ ‡‡‡‡‡‡‡‡‡‡‡— ‰ j                   ‰   \  Š}Š}	}‰ j                  t        j                  j                  ‰‰‰‰‰‰‰g‰«	      ry ‰ j                  ‰ ‰‰‰«      Št        j                  j                  «       }
|
�It        «       Š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@   rA   rž   s   €€€€€rC   Úasync_compilez.JITFunction._do_compile.<locals>.async_compilei  s!   ø€ Ø—|‘| C°À×@PÑ@PÐ\d�|ÓeÐerD   c                ót   •
— | ‰‰<   ‰j                  t        j                  j                  ‰‰‰	‰‰‰‰g‰
«	       y rF   )r�  r   r­  Újit_post_compile_hook)r½  r  r¦  r{  rº  r  r   r@   rû   rž   rë   s    €€€€€€€€€€rC   Úfinalize_compilez1JITFunction._do_compile.<locals>.finalize_compilel  s;   ø€ Ø$*�˜SÑ!Ø—‘¤§¡× CÑ CÀSÈ)ÐU[Ð]cÐeoØ '¨%¨°&õ:rD   )rž   r   )r±  r�  r   r­  Újit_cache_hookr“  r   Úactive_moderw   r   r   Úsubmitr†  r  rß  )r@   r  rû   r{  r¦  r   r  rë   r`   r  Ú
async_moder\   rÝ  rà  r½  rÜ  rº  rA   rž   s   ````````       @@@@rC   r³  zJITFunction._do_compile[  s0  ÿû€ Ø.2×.@Ñ.@ÀÑ.HÑ+ˆ�a˜ ¨!à�?‰?œ5Ÿ=™=×7Ñ7¸¸iÈÐQWÐYcÐelÐotÐnuØ!ô#àØ�n‰n˜T 9¨j¸%Ó@ˆä#×/Ñ/×3Ñ3Ó5ˆ
ØÐ!ä6Ó8ˆHÜ% c¨7°G¸XÓFˆI÷fð f÷:õ :ð
  ×&Ñ& y°-ÐAQÓRˆFØ &ˆL˜Ñð ˆð	 —\‘\ #¨f¸g×>NÑ>N�\ÓOˆFØ &ˆL˜ÑØ�O‰OœEŸM™M×?Ñ?ÀÀiÐQWÐY_ÐakÐmtØ"˜G Vô-àˆrD   c                ó   — t        d«      ‚)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rY   ©r@   rŠ   rï   s      rC   Ú__call__zJITFunction.__call__z  s   € ÜÐWÓXÐXrD   c                óP   — d| j                   › d| j                  j                  › d�S )NzJITFunction(rú   r  )rR   r  r­   rH   s    rC   Ú__repr__zJITFunction.__repr__}  s&   € Ø˜dŸk™k˜]¨!¨D¯G©G×,@Ñ,@Ð+AÀÐCÐCrD   )r«   zbool | None)NNNNNNN)rZ   rK   r­   rx  r�  r�  r™  r©  rç   r…  r2   rØ  r³  rç  ré  r°   r±   s   @rC   rø   rø   b  sZ   ø„ òð.
ð 
ó.
ò`(ò/ò5ò07òrFð mqØ;?õ$ òL3
òjò>YöDrD   rø   c                 ó   — y rF   rà   r*  s    rC   Újitrë  †  s   € àrD   ©rÈ  r…  rµ  rÉ   rÊ   r«  rV   c                 ó   — y rF   rà   rì  s          rC   rë  rë  ‹  s   € ð rD   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«  rV   r…  rµ  )rn   r   r­  Ú	interpretÚinterpreterrð  rø   )	r  rð  r«  rÉ   rÊ   rµ  rV   r…  rÈ  s	     €€€€€€€rC   Ú	decoratorzjit.<locals>.decorator¶  sj   ø€ Ü˜Œ|Ðˆ|Ü�=‰=×"Ò"Ý8Ù& r°7ÐN_ØFdÐlqØ08¸tÐUdôfð fô ØØØ"3Ø/MØØ!ØØ /ô	ð 	rD   ©r  r    r«   zJITFunction[T]rà   )	r  rÈ  r…  rµ  rÉ   rÊ   r«  rV   ró  s	    ``````` rC   rë  rë  ™  s&   þ€ ÷:ó ð& 
€~Ù˜‹}Ðð ÐrD   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)rB   rZ   rK   rí   )r‹   s    rC   rî   zMockTensor.wrap_dtypeÛ  s.   € à�=‰=×!Ñ! WÒ,°·±À7Ò1JÜ˜c“?Ð"Øˆ
rD   Nc                ó*   — |€dg}|| _         || _        y )Nr   )rÁ   Úshape)r@   rÁ   rù  s      rC   r2   zMockTensor.__init__á  s   € Øˆ=Ø�CˆEØˆŒ
Øˆ�
rD   c                ó†   — dg}| j                   dd  D ]  }|j                  |d   |z  «       Œ t        t        |«      «      S )Nr   r¶   )rù  r  rl  Úreversed)r@   ÚstridesÚsizes      rC   ÚstridezMockTensor.strideç  sG   € Ø�#ˆØ—J‘J˜q˜r�Nò 	/ˆDØ�N‰N˜7 2™;¨Ñ-Õ.ð	/ä”X˜gÓ&Ó'Ð'rD   c                  ó   — y©Nr   rà   rà   rD   rC   Údata_ptrzMockTensor.data_ptrí  ó   € àrD   c                  ó   — yr   rà   rà   rD   rC   Ú	ptr_rangezMockTensor.ptr_rangeñ  r  rD   rF   )
rZ   rK   r­   r®   Ústaticmethodrî   r2   rþ  r  r  rà   rD   rC   rí   rí   Õ  sM   „ ñð
 ñó ðó
ò(ð ñó ðð ñó ñrD   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 rF   )rÁ   ÚbaseÚdatar{  rù  )r@   r	  rÁ   s      rC   r2   zTensorWrapper.__init__ø  s5   € ØˆŒ
ØˆŒ	Ø—I‘IˆŒ	Ø—k‘kˆŒØ—Y‘Y—_‘_ˆ�
rD   c                ó6   — | j                   j                  «       S rF   )r	  r  rH   s    rC   r  zTensorWrapper.data_ptrÿ  s   € Ø�y‰y×!Ñ!Ó#Ð#rD   c                ó4   —  | j                   j                  |Ž S rF   )r	  rþ  )r@   rŠ   s     rC   rþ  zTensorWrapper.stride  s   € Øˆt�y‰y×Ñ Ð&Ð&rD   c                ó<   — d| j                   › d| j                  › d�S )NzTensorWrapper[rz  r  )rÁ   r	  rH   s    rC   Ú__str__zTensorWrapper.__str__  s   € Ø §
¡
˜|¨2¨d¯i©i¨[¸Ð:Ð:rD   c                ó6   — | j                   j                  «       S rF   )r	  Úelement_sizerH   s    rC   r  zTensorWrapper.element_size  s   € Ø�y‰y×%Ñ%Ó'Ð'rD   c                ó^   — t        | j                  j                  «       | j                  «      S rF   )r  r	  ÚcpurÁ   rH   s    rC   r  zTensorWrapper.cpu  s   € Ü˜TŸY™YŸ]™]›_¨d¯j©jÓ9Ð9rD   c                óN   — | j                   j                  |j                   «       y rF   )r	  Úcopy_)r@   Úothers     rC   r  zTensorWrapper.copy_  s   € Ø�	‰	�‰˜Ÿ
™
Õ#rD   c                ó^   — t        | j                  j                  «       | j                  «      S rF   )r  r	  ÚclonerÁ   rH   s    rC   r  zTensorWrapper.clone  s   € Ü˜TŸY™YŸ_™_Ó.°·
±
Ó;Ð;rD   c                ó`   — t        | j                  j                  |«      | j                  «      S rF   )r  r	  ÚtorÁ   )r@   r{  s     rC   r  zTensorWrapper.to  s    € Ü˜TŸY™YŸ\™\¨&Ó1°4·:±:Ó>Ð>rD   c                ó`   — t        | j                  j                  |«      | j                  «      S rF   )r  r	  Ú	new_emptyrÁ   )r@   Úsizess     rC   r  zTensorWrapper.new_empty  s"   € Ü˜TŸY™Y×0Ñ0°Ó7¸¿¹ÓDÐDrD   Nrß   )rZ   rK   r­   r2   r  rþ  r  r  r  r  r  r  r  rà   rD   rC   r  r  ö  s5   „ ò%ò$ò'ó;ò(ò:ò$ò<ò?óErD   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Á   Úhasattrr¢   rk   )ÚtensorrÁ   s     rC   Úreinterpretr     sk   € Ü�&œ-Ô(Ø�F—K‘K×%Ñ%Ò%à—;‘;Ðô ! §¡¨eÓ4Ð4Ü	�˜Ô	$ä˜V UÓ+Ð+äÐ/´°V³¨~¸QÐ?Ó@Ð@rD   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.   rW   r  rC  Úco_filenamer0  rq  r/  r»   rP   )r  Úbase_fnÚ	file_nameÚ
begin_lineÚidxÚlines         rC   Úget_jit_fn_file_liner(  *  s�   € Ø€GÜ˜¤+Ô.Ø—*‘*ˆô ˜¤+Õ.à—
‘
×#Ñ#×/Ñ/€IØ×-Ñ-€Jô ˜wŸ™Ó/ò ‰	ˆˆTØ�:‰:‹<×"Ñ" 6Õ*Ø˜#ÑˆJØØ�jÐ Ð ð	ð �jÐ Ð rD   c                  ó(   — e Zd Zd„ Zed„ «       Zd„ Zy)ÚBoundConstexprFunctionc                ó    — || _         || _        y rF   )Ú__self__Ú__func__)r@   Úinstancer  s      rC   r2   zBoundConstexprFunction.__init__>  s   € Ø ˆŒØˆ�rD   c                ó.   — | j                   j                  S rF   )r-  r\   rH   s    rC   r\   z BoundConstexprFunction.cache_keyB  s   € à�}‰}×&Ñ&Ð&rD   c                óB   —  | j                   | j                  g|¢­i |¤ŽS rF   )r-  r,  ræ  s      rC   rç  zBoundConstexprFunction.__call__F  s!   € Øˆt�}‰}˜TŸ]™]Ð<¨TÒ<°VÑ<Ð<rD   N)rZ   rK   r­   r2   r¯   r\   rç  rà   rD   rC   r*  r*  <  s    „ òð ñ'ó ð'ó=rD   r*  c                  ó0   ‡ — e Zd Zˆ fd„Zd„ Zddœd„Zˆ xZS )ÚConstexprFunctionc                ó$   •— t         ‰| �  |«       y rF   )r1   r2   )r@   r  rB   s     €rC   r2   zConstexprFunction.__init__L  s   ø€ Ü‰Ñ˜ÕrD   c                ó"   — |�t        || «      S | S rF   )r*  )r@   r
  Úobjclasss      rC   Ú__get__zConstexprFunction.__get__O  s   € àˆ?Ü)¨#¨tÓ4Ð4ØˆrD   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_constexprrg   )r¸   r9  rg   r  r  r   r­  rñ  )
r@   r7  rŠ   rï   r9  rg   r	  r^   r¥  Úress
             rC   rç  zConstexprFunction.__call__U  sŽ   € ßHà15Ö6¨AÑ$ QÕ'Ð6ˆÐ6Ø;A¿<¹<»>×J±°!°Q�!Ñ)¨!Ó,Ñ,ÐJˆÑJð ˆd�g‰g�tÐ&˜vÑ&ˆàÐàˆJô �=‰=×"Ò"ØˆJÙ˜‹~Ðùò 7ùÛJs
   �B±B)rZ   rK   r­   r2   r6  rç  r°   r±   s   @rC   r2  r2  J  s   ø„ ôòð )-÷ ð rD   r2  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.
    )r2  r*  s    rC   Úconstexpr_functionr<  h  s   € ô ˜RÓ Ð rD   rß   )Frô  )r…  úOptional[Callable]rµ  r=  rÉ   úOptional[Iterable[int | str]]rÊ   r>  r«  úOptional[bool]rV   r?  r«   zCallable[[T], JITFunction[T]]rF   )r  zOptional[T]r…  r=  rµ  r=  rÉ   r>  rÊ   r>  r«  r?  rV   r?  r«   zKernelInterface[T])HÚ
__future__r   r   rz   r/   r4   rM   r“   r4  r9  r7  Úcollectionsr   Údataclassesr   Ú	functoolsr   Útypingr   r	   r
   r   r   r   r   r   r   Útriton.backendsr   Útypesr   rL   r   r   r   Ú_utilsr   r   r   r   Úcacher   Útriton._C.libtritonr   r   r   r<   r;   r    ÚNodeVisitorr"   r½   rÅ   rä   ræ   r  r'  r+  rW   rÉ  re  ro  rr  rø   rë  rí   r  r   r(  r*  r2  r<  rà   rD   rC   ú<module>rK     s;  ðß ,Û 
Û Û Û Û Û Û 	Û Ý #Ý !Ý %ß ]× ]Õ ]å 'Ý Ý Ý Ý ß `Ó `Ý  ß [Ñ [à!€Ø3€áˆCƒL€ôg!˜Ÿ™ô g!ó^D÷4/>ñ />ódTôZ�g˜a‘jô Zò&ò$9*òx0÷i1ñ i1ðX !#Ð ð ÷ð ó ðòò0	ô\D�+˜¨qÑ1ô \DðH	 
òó 
ðð 
ð Ø#Ø*.Ø7;ØDHØ Ø#ñ
ð ð
ð (ð	
ð
 5ð
ð %Bð
ð ð
ð ð
ð #ò
ó 
ð
ð ð4ð Ø#Ø*.Ø7;ØDHØ Ø#ñ4Øð4ð ð	4ð
 (ð4ð 5ð4ð %Bð4ð ð4ð ð4ð ó4÷xñ ÷B"Eñ "EòJAò!ô$=˜[ô =ô˜ô ó<!rD   