Ë
    ©š¾jË®  ã            	      ó„  — d dl mZmZ d dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
Z
d dlmZ d dlmZ d dlmZ d dlmZmZmZmZmZmZmZmZmZmZmZmZ d dlmZ d dl m!Z! d	d
l"m#Z# ddl$m$Z$ ddl"m%Z% d	dl&m'Z'm(Z(m)Z)m*Z* ddl+m,Z, d dl-m.Z.m/Z/m0Z0 dZ1dZ2 e	jf                  de	jh                  «      Z5 ed«      Z6 ed«      Z7 ed«      Z8 ed«      Z9 G d„ dejt                  «      Z;d;d„Z< G d„ d«      Z=d<d„Z> G d„ dee6   «      Z?d „ Z@d!„ ZAd"„ ZB G d#„ d$«      ZCi ZDe G d%„ d&«      «       ZEd'„ ZFd(„ ZG G d)„ d*eCe?e6   «      ZHed=d+„«       ZIedddddddd,œ	 	 	 	 	 	 	 	 	 	 	 	 	 d>d-„«       ZI	 d?dddddddd,œ	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d@d.„ZI G d/„ d0«      ZJ G d1„ d2«      ZKd3„ ZLd4„ ZMd5„ ZN G d6„ d7eC«      ZO G d8„ d9eCee6   «      ZPdAd:„ZQy)Bé    )ÚannotationsÚdivisionN)Údefaultdict)Ú	dataclass)Úcached_property)ÚCallableÚConcatenateÚGenericÚIterableÚOptionalÚ	ParamSpecÚTYPE_CHECKINGÚTypeVarÚoverloadÚDictÚAnyÚTuple)ÚBaseBackend)Ú
ModuleTypeé   )Úknobsé   )Údriver)Ú_async_compile)Úfind_paths_ifÚget_iterable_pathÚtype_canonicalisation_dictÚis_namedtuple)Úget_cache_key)Úget_cache_invalidating_env_varsÚnative_specialize_implÚirztriton.languagez"triton.experimental.gluon.languagez!^(?P<indent>[ \t]*)def\s+\w+\s*\(ÚTÚPÚRÚUc                  óx   ‡ — e Zd ZdZdˆ fd„Zed„ «       Zd„ Zdd„Zd„ Z	d„ Z
d„ Zd	„ Zd
„ Zd„ Zd„ Zd„ Zd„ Zˆ xZS )ÚDependenciesFinderaÑ  
    This AST visitor is used to find dependencies of a JITFunction. This can
    be used to invalidate a JITFunction's hash when its source code -- or
    that of its dependencies -- changes.

    This visitor also keeps track of the global variables touched by the
    JITFunction.  When we launch the kernel, we check that these have the same
    values as they did when we ran this visitor.  If not, we raise an error (or
    otherwise we could recompile).
    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)Úselfr9   r>   r?   ÚsrcÚ	__class__s        €ú\/home/spazio/kingdom_calls/work5/transenv/lib/python3.12/site-packages/triton/runtime/jit.pyr8   zDependenciesFinder.__init__3   st   ø€ Ü‰ÑÔØˆŒ	Ü—n‘n S§Z¡Z°Ó%8Ó9ˆŒð ˆŒØ"ˆŒò*
ˆÔ&ô ÜØØð	"
ˆÔð" TVˆÔà*/ˆÕ'ó    c                ó6   — | j                   j                  «       S ©N)r=   Ú	hexdigest©rF   s    rI   ÚretzDependenciesFinder.ret^   s   € à�{‰{×$Ñ$Ó&Ð&rJ   c                ó>  — t        |t        «      sJ ‚|j                  }| j                  j	                  «       |j                  j	                  «       z  D ]_  }|\  }}| j                  |   \  }}|j                  |   \  }}||k7  sŒ2t        d|› d|› d| j                  › d|j                  › d|› d�«      ‚ | 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*   )r4   ÚJITCallableÚ	cache_keyrD   ÚkeysÚRuntimeErrorr9   Ú__name__ÚupdateÚstrr3   r=   r<   )rF   ÚfuncÚfunc_keyÚkÚvar_nameÚ_Úv1Úv2s           rI   Ú_update_hashzDependenciesFinder._update_hashb   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Õ4rJ   c                óz  — ddl m} |�t        |«      t        u ry t	        |dd«      r`| j
                  j                  t        |j                  «      j                  d«      «       |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__Fr*   Ú__triton_builtin__Ú
__module__Ú ztriton.language.extra.libdevicez!Unsupported function referenced: )Úlanguage.corerd   Útyper   r3   r=   rX   rY   Ú__annotations__r<   Ú
hash_attrsÚrecord_referencer4   rS   ra   ÚcallablerV   rE   r5   ÚdeepcopyrD   Úid)rF   ÚvalÚvar_dictr9   rd   Úattrs         rI   rm   z#DependenciesFinder.record_referencet   s  € Ý-ð ˆ;œ$˜s›)¤zÑ1Øä�3Ð.°Ô6Ø�K‰K×Ñœs 3×#6Ñ#6Ó7×>Ñ>¸wÓGÔHØŸ™ò ,�Ø×%Ñ% 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ØrJ   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)r>   Úgetr?   )r9   rq   rF   s     €rI   Úname_lookupz2DependenciesFinder.visit_Name.<locals>.name_lookup¤   sZ   ø€ Ø—,‘,×"Ñ" 4¨Ó.ˆCØˆØ˜DŸL™LÐ(Ð(Ø—.‘.×$Ñ$ T¨4Ó0ˆCØˆØ˜DŸN™NÐ*Ð*ØrJ   )rj   ÚctxÚastÚStorerp   Úlocal_namesr@   rm   )rF   Únoderx   rq   rr   s   `    rI   Ú
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Øˆ
rJ   c                ó^   — |j                   D �cg c]  }| j                  |«      ‘Œ c}S c c}w rL   )ÚeltsÚvisit)rF   r}   Úelts      rI   Úvisit_TuplezDependenciesFinder.visit_Tuple´   s$   € ð ,0¯9©9Ö5 C�—
‘
˜3•Ò5Ð5ùÒ5s   �*c                óh  — | j                  |j                  «      }t        |t        j                  «      r6| j                  |j                  «      }t        |t        j                  «      rŒ6t        |dd«      }|�|| j                  v ry t        ||j                  d «      }| j                  |«       |S )NrW   rh   )	r�   Úvaluer4   rz   Ú	Attributer3   rC   rs   rm   )rF   r}   ÚlhsÚlhs_namerO   s        rI   Ú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 dÓ+ˆØ×Ñ˜cÔ"Øˆ
rJ   c                ó–   — |j                   j                   D �ch c]  }|j                  ’Œ c}| _        | j                  |«       y c c}w rL   )ÚargsÚargr|   Úgeneric_visit)rF   r}   rŒ   s      rI   Ú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)rE   r�   )ÚdefaultsÚexprrF   s     €rI   Úvisit_defaultsz:DependenciesFinder.visit_arguments.<locals>.visit_defaultsÕ   sS   ø€ ð8Ø×:Ò:Ð:Ð:Ø26�Ô/Ø$ò )�DØÑ'ØŸ
™
 4Õ(ñ)ð 38�Õ/ø°%�Õ/ús   ƒ<  < ¼	A)
Ú	itertoolsÚchainÚposonlyargsr‹   ÚvarargÚ
kwonlyargsr�   Úkw_defaultsÚkwargr‘   )rF   r}   r“   rŒ   s   `   rI   Ú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—}‘}Õ%rJ   c                óF   ‡ ‡— ‰ j                  |«      }ˆˆ fd„Š ‰|«       y )Nc                ó|   •— t        | t        «      r| D ]
  } ‰|«       Œ y ‰j                  j                  | «       y rL   )r4   r/   r|   Úadd)ÚtÚsubÚ_addrF   s     €€rI   r¡   z0DependenciesFinder.visitAssnTarget.<locals>._addí   s9   ø€ Ü˜!œTÔ"Øò �CÙ˜•Iñð × Ñ ×$Ñ$ QÕ'rJ   )r�   )rF   r}   Útargetr¡   s   `  @rI   ÚvisitAssnTargetz"DependenciesFinder.visitAssnTargeté   s   ù€ à—‘˜DÓ!ˆõ	(ñ 	ˆV�rJ   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�   ©rF   r}   s     rI   Úvisit_AssignzDependenciesFinder.visit_Assignö   sG   € Üˆt�|‰|Ó Ò!ô
 ÐPÓQÐQà×Ñ˜TŸ\™\¨!™_Ô-ð 	×Ñ˜4Õ rJ   c                ó\   — | j                  |j                  «       | j                  |«       y rL   ©r£   r¢   r�   r§   s     rI   Úvisit_AnnAssignz"DependenciesFinder.visit_AnnAssign  ó$   € Ø×Ñ˜TŸ[™[Ô)ð 	×Ñ˜4Õ rJ   c                ó\   — | j                  |j                  «       | j                  |«       y rL   rª   r§   s     rI   Ú	visit_ForzDependenciesFinder.visit_For	  r¬   rJ   )ÚreturnÚNonerv   )rW   rg   Ú__qualname__Ú__doc__r8   ÚpropertyrO   ra   rm   r~   rƒ   r‰   rŽ   r›   r£   r¨   r«   r®   Ú__classcell__©rH   s   @rI   r(   r(   '   s[   ø„ ñ	õ)0ðV ñ'ó ð'ò5ó$&òPò06ò
	ò!ò
&ò@ò!ò!ö!rJ   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.Ú_trh   )Útriton.language.coreÚlanguageÚcorer4   rY   ÚstripÚ
startswithÚremoveprefixÚ_normalize_tyÚendswithÚpointer_typeÚ
element_tyÚdtyper9   rj   rW   r   rw   Úreplace)Útyr¾   s     rI   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ÐCrJ   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 rL   )ÚnumÚ_paramÚdo_not_specializeÚdo_not_specialize_on_alignment)rF   rÌ   ÚparamrÎ   rÏ   s        rI   r8   zKernelParam.__init__2  s    € àˆŒØˆŒØ!2ˆÔØ.LˆÕ+rJ   c                ó.   — | j                   j                  S rL   )rÍ   r9   rN   s    rI   r9   zKernelParam.name9  s   € à�{‰{×ÑÐrJ   c                óÐ   — | j                   j                  r1| j                   j                  t        j                  j                  k(  ryt        | j                   j                  «      S )Nrh   )rÍ   Ú
annotationÚinspectÚ	ParameterÚemptyrÂ   rN   s    rI   rÓ   zKernelParam.annotation=  sD   € à�{‰{×%Ò%¨¯©×)?Ñ)?Ä7×CTÑCT×CZÑCZÒ)ZØÜ˜TŸ[™[×3Ñ3Ó4Ð4rJ   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   rh   )rÓ   rÀ   Úsetr   Úvalues)rF   Úas     rI   Úannotation_typezKernelParam.annotation_typeC  s]   € à�O‰OˆØ�<‰<˜ÔØ�!�"�‰AØ�\‰\˜#ÔØ�!�"�ˆAØ”Ô.×5Ñ5Ó7Ó8Ñ8Ø—?‘?Ð"ØrJ   c                ó   — d| j                   v S ©Nrd   )rÓ   rN   s    rI   Úis_constexprzKernelParam.is_constexprN  s   € à˜dŸo™oÐ-Ð-rJ   c                ór   — | j                   ryd| j                  v xs | j                  j                  d«      S )NFr·   r¹   )rÞ   rÓ   rÀ   rN   s    rI   Úis_constzKernelParam.is_constR  s1   € à×ÒØØ˜$Ÿ/™/Ð)ÒM¨T¯_©_×-GÑ-GÈÓ-MÐMrJ   c                ó.   — | j                   j                  S rL   )rÍ   ÚdefaultrN   s    rI   râ   zKernelParam.defaultX  s   € à�{‰{×"Ñ"Ð"rJ   c                ód   — | j                   j                  t        j                  j                  k7  S rL   )rÍ   râ   rÔ   rÕ   rÖ   rN   s    rI   Úhas_defaultzKernelParam.has_default\  s#   € à�{‰{×"Ñ"¤g×&7Ñ&7×&=Ñ&=Ñ=Ð=rJ   N)rÌ   r+   rÐ   zinspect.ParameterrÎ   ÚboolrÏ   rå   ©r¯   rY   )rW   rg   r±   r²   r8   r   r9   rÓ   rÛ   rÞ   rà   r³   râ   rä   © rJ   rI   rÊ   rÊ   /  s¢   „ ÙLðMØ15óMð ñ ó ð ð ò5ó ð5ð
 òó ðð ñ.ó ð.ð ñNó ðNð
 ñ#ó ð#ð ñ>ó ñ>rJ   rÊ   c                ó6   — d}d}t        t        | |||«      d   S )NFTr   )r!   r   )rŒ   Ú
specializerà   Úaligns       rI   Úmangle_typerë   a  s%   € Ø€HØ€EÜ!¤+¨s°H¸jÈ%ÓPÐQRÑSÐSrJ   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)rF   rñ   r‹   Úkwargss       rI   rò   zKernelInterface.warmupj  s*   € Øˆt�x‰x´´Z×5JÑ5JÈDÓ1Q˜T¨$Ñ\ÐU[Ñ\Ð\rJ   c               ó   — t        d«      ‚)Nzrun not implemented)ÚNotImplementedError)rF   rñ   rò   r‹   rö   s        rI   rî   zKernelInterface.runm  s   € Ü!Ð"7Ó8Ð8rJ   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ñ   rF   s     €€rI   ú<lambda>z-KernelInterface.__getitem__.<locals>.<lambda>v  s   ø€  x t§x¡xÈ$°TÀ%Ñ'YÐRXÑ'Y€ rJ   rç   )rF   rñ   s   ``rI   Ú__getitem__zKernelInterface.__getitem__p  s   ù€ ô ZÐYrJ   N)r¯   r#   )rW   rg   r±   rk   rò   rî   rü   rç   rJ   rI   rí   rí   g  s   … Ø	
ƒFò]ò9ôZrJ   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Æ   rd   ÚJITFunctionÚjit_functionú:r   )	r9   Ú	signatureÚconstant_keysÚconstant_valsÚ
attrs_keysÚ
attrs_valsÚoptionsÚkeyr¢   )ÚitemsrH   rW   rY   r…   ÚmoduleÚfnr±   ÚjsonrU   r/   rÙ   Ú__dict__Údumps)r9   r  Ú	constantsÚattrsr  r  r¢   r…   r  ÚxÚobjÚserialized_objs               rI   Úserialize_specialization_datar  z  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                óÂ   — | \  }}|j                   t        j                  j                  k(  rd|› �S |j                  t        j                  j
                  u r|S |› d|› �S )Nr¸   z	=default_)ÚkindrÔ   rÕ   ÚVAR_POSITIONALrâ   rÖ   )Ú
name_paramr9   rÐ   s      rI   rŒ   z+create_function_from_signature.<locals>.arg­  s^   € Ø ‰ˆˆeØ�:‰:œ×*Ñ*×9Ñ9Ò9Ø�t�f�:ÐØ—}‘}¬×(9Ñ(9×(?Ñ(?Ñ?ˆtÐ]ÈÀvÈYÐW[ÐV\ÐE]Ð]rJ   z
def dynamic_func(z	**optionsz):
    params = {ú'z': z}
    specialization = [ú,z-]
    return params, specialization, options
Údefault_Úspecialize_implÚbackendrS   Údynamic_func)r,   Ú
parametersÚziprU   rÞ   Úappendrà   rÎ   rÏ   rÛ   r4   rY   Újoinr/   ró   r  râ   rÔ   rÕ   rÖ   r!   rS   Úexec)ÚsigÚkparamsr$  Úspecializationr9   Úkprà   ré   rê   rO   rŒ   Ú	func_bodyrÐ   Úfunc_namespacer#  s                  rI   Úcreate_function_from_signaturer1  Œ  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ò,^ðØ—)‘)œ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ú.)rg   r±   ©r
  s    rI   Úget_full_namer5  Í  s   € Ø�m‰mˆ_˜A˜bŸo™oÐ.Ð/Ð/rJ   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)rS   c                óX  — || _         t        j                  |«      | _        	 t        j                  |«      \  | _        | _        t        |«      | _	        t        j                  «       | _        dj                  | j                  «      }|j                  j                  | _        t#        | j                  | j
                  «      | _        t'        |«      | _        t+        j,                  |«      }|t/        j0                  d|t.        j2                  «      j5                  «       d  }|| _        d | _        i | _        |j<                  | _        |j>                  | _        |j@                  | _         |jB                  | _!        |jD                  | _"        y # t        $ r}t        d«      |‚d }~ww xY w)Nz1@jit functions should be defined in a Python filerh   z^def\s+\w+\s*\()#r
  rÔ   r  ÚgetsourcelinesÚraw_srcÚstarting_line_numberÚOSErrorÚ
ValueErrorr5  Ú_fn_nameÚ	threadingÚRLockÚ
_hash_lockr)  Ú__code__Úco_filenameÚ	file_nameÚget_def_line_numberÚdef_file_line_numberÚget_def_col_numberÚdef_file_col_numberÚtextwrapÚdedentÚreÚsearchÚ	MULTILINEÚstartÚ_srcÚhashrD   r²   rW   r±   Ú__globals__rg   )rF   r
  ÚeÚraw_src_strrG   s        rI   r8   zJITCallable.__init__Ó  sS  € ØˆŒÜ ×*Ñ*¨2Ó.ˆŒð	YÜ6=×6LÑ6LÈRÓ6PÑ3ˆDŒL˜$Ô3ô & bÓ)ˆŒÜ#Ÿ/™/Ó+ˆŒð —g‘g˜dŸl™lÓ+ˆð Ÿ™×0Ñ0ˆŒÜ$7¸¿¹Àd×F_ÑF_Ó$`ˆÔ!Ü#5°kÓ#BˆÔ ä�o‰o˜kÓ*ˆØ”"—)‘)Ð.°´R·\±\ÓB×HÑHÓJÐKÐLˆØˆŒ	ØˆŒ	ð TVˆÔð —z‘zˆŒØŸ™ˆŒØŸO™OˆÔØŸ>™>ˆÔØŸ-™-ˆ�øôE ò 	YÜÐPÓQÐWXÐXûð	Yús   £"F Æ	F)ÆF$Æ$F)c                ó  — | j                   }|j                  €| j                  S t        |j                  j
                  |j                  «      D ��ci c]  \  }}||j                  “Œ }}}| j                  |z  S c c}}w rL   )r
  Ú__closure__rP  r'  rA  Úco_freevarsÚcell_contents)rF   r
  r9   Úcellr?   s        rI   Ú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} | j"                  j%                  «       D ���cg c]  \  \  }}\  }}t+        ||«      r||f‘Œ }}}}|j-                  d„ ¬«       | xj                  t        |«      z  c_        t/        j0                  | j                  j3                  d«      «      j5                  «       | _        d d d «       | j                  S c c}}}w # 1 sw Y   | j                  S xY w)Nz
recursion:)r9   r>   r?   rG   r   rc   c                ó(   — | d   t        | d   «      fS )Nr   r   )Úrepr)Úitems    rI   rû   z'JITCallable.cache_key.<locals>.<lambda>  s   € °T¸!±W¼dÀ4ÈÁ7»mÐ4L€ rJ   )r  r*   )r@  rO  r=  rÔ   Úgetclosurevarsr
  r?   r(   rP  rG   r�   ÚparserO   rY   r:  ÚdictÚsortedrD   r  r¼   rd   r4   Úsortr:   r;   r<   rM   )rF   r?   Údependencies_finderrd   r9   r^   rq   Úconstexpr_globalss           rI   rT   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à<@×<QÑ<Q×<WÑ<WÓ<Y÷!@ð !@Ù%8¡Y d¨A±°°aÜ$.¨s°IÔ$>ð #'¨¢ð !@Ðò !@ð ×"Ñ"Ñ'LÐ"ÔMØ�IŠIœÐ.Ó/Ñ/�IÜŸ™ t§y¡y×'7Ñ'7¸Ó'@ÓA×KÑKÓMˆDŒI÷'	Nð( �y‰yÐùô!@÷	Nð( �y‰yÐús$   �G%¯DG%Ä4 G
ÅA6G%ÇG%Ç%G9c                ó,   — t        | j                  «      S rL   )rO  rT   rN   s    rI   Ú__hash__zJITCallable.__hash__  s   € Ü�D—N‘NÓ#Ð#rJ   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   r^  rN  r4   ÚModuler,   ÚbodyÚFunctionDef)rF   Útrees     rI   r^  zJITCallable.parse"  s_   € Ü�y‰y˜Ÿ™Ó#ˆÜ˜$¤§
¡
Ô+Ð+Ð+Ü�4—9‘9‹~ Ò"Ð"Ð"Ü˜$Ÿ)™) A™,¬¯©Ô8Ð8Ð8ØˆrJ   c                ó   — ddl m}  || «      S )Nr   )Úconstexpr_type)r¼   rl  )rF   rl  s     rI   rj   zJITCallable.type)  s   € å7Ù˜dÓ#Ð#rJ   c                 ó   — y rL   rç   )rF   Úhandless     rI   Ú_flatten_irzJITCallable._flatten_ir.  s   € ØrJ   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)rO  rN  )rF   Únew_srcs     rI   Ú_unsafe_update_srczJITCallable._unsafe_update_src1  s   € ð ˆŒ	Øˆ�	rJ   c                ó   — t        d«      ‚)NzqCannot set attribute 'src' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)ÚAttributeErrorrN   s    rI   Ú_set_srczJITCallable._set_src;  s   € Üð (ó )ð 	)rJ   c                ó   — | j                   S rL   )rN  rN   s    rI   Ú_get_srczJITCallable._get_src@  s   € Ø�y‰yÐrJ   )ÚfgetÚfsetNræ   )rn  zlist[ir.value]r¯   r°   )rW   rg   r±   r8   rX  r³   rT   re  r^  rj   ro  rr  ru  rw  rG   rç   rJ   rI   rS   rS   Ñ  sa   „ ò'(òR,ð òó ðò0$òð ñ$ó ð$óòò)ò
ñ ˜ xÔ
0�CrJ   rS   c                  ó,   — e Zd ZU ded<   ded<   ded<   y)ÚJitFunctionInfor   r	  rY   r9   rþ   rÿ   N)rW   rg   r±   rk   rç   rJ   rI   r{  r{  I  s   … àÓØ
ƒIØÔrJ   r{  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rL   rç   )Ú.0rŒ   Úreplace_callabless     €rI   ú	<genexpr>z?compute_cache_key.<locals>.replace_callables.<locals>.<genexpr>^  s   øè ø€ Ò?°CÑ*¨3×/Ñ?ùs   ƒ)r4   r/   r   rH   ÚtuplerS   rT   )r  rŒ   Úresultsr€  s      €rI   r€  z,compute_cache_key.<locals>.replace_callablesW  s’   ø€ Ü�cœ4Ô Ø69Ö:¨sÑ% cÕ*Ò:Ð:Ü˜3ÔØ9<Ö=°#Ñ(¨Õ-Ð=ˆGÐ=Ø �3—=‘= 'Ð*Ð*Ü˜œUÔ#ÜÓ?¸3Ô?Ó?Ð?Ü˜œ[Ô)Ø—=‘=Ð Øˆ
ùò ;ùâ=s
   –B·B)r‚  rY   rw   )Úkernel_key_cacher-  r  r  rT   r€  s        @rI   Úcompute_cache_keyr…  P  se   ø€ Ü�Ó ¤# g£,Ð
/€CØ ×$Ñ$ S¨$Ó/€IØÐØÐô
ô Ñ% nÓ5Ó6¼¸W»ÑE€IØ%Ð�SÑØÐrJ   c                ó~   — t        | t        «      s| S t        | «      D ]  \  }}t        |«      | |<   Œ t	        | «      S rL   )r4   r/   Ú	enumerateÚconvert_to_tuple_if_listr‚  )r\  ÚiÚnested_values      rI   rˆ  rˆ  h  sD   € ä�dœDÔ!Øˆô % T›?ò 9‰ˆˆ<Ü*¨<Ó8ˆˆQŠð9ô �‹;ÐrJ   c                  óª   ‡ — e Zd Zd„ Z	 	 dd„Zd„ Zd„ Zd„ Zd„ Zd„ Z		 	 dˆ fd„	Z
d	„ Zd
„ Zdd„Zereddd„«       Ze	 d	 	 	 dd„«       Zdd„Zd„ Zˆ xZS )rþ   c                 ó   — y)NFrç   rN   s    rI   Úis_gluonzJITFunction.is_gluonv  s   € ØrJ   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  r[  r
  ÚcompileÚis_manual_warmupÚalready_compiled)r
  r±   rg   r)  r'  Úparamsr9   r‘  r’  r“  r”  r•  r5  r  r–  r{  )rF   Úhookr  r  r¢   r�  r  r  r—  r™  r9   r	  rÐ   rÈ   Ú	arg_reprsr[  Ú	full_namer˜  rö   s                      rI   Ú
_call_hookzJITFunction._call_hooky  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(  )rF   rž  s     rI   Úadd_pre_run_hookzJITFunction.add_pre_run_hook©  s$   € ô
 ˜Œ~Ðˆ~Ø×Ñ×!Ñ! $Õ'rJ   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_targetr1  r  r�  )rF   r¦  rš  r§  r¨  r¢   r$  Úbinders           rI   Úcreate_binderzJITFunction.create_binder±  sd   € ÷ 	PÓOÜ—‘×1Ñ1Ó3ˆÙ˜vÓ&ˆØ,ˆÔØˆŒØ"ˆŒÜ/°·±ÀÇÁÈWÓUˆØ�2�v˜w¨Ð.Ð.rJ   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     rI   rû   z(JITFunction._pack_args.<locals>.<lambda>Í  s   € ¸3À+Ñ;M€ rJ   rd   rh   r   c                ó"   — t        |t        «      S rL   )r4   rY   )r^   r  s     rI   rû   z(JITFunction._pack_args.<locals>.<lambda>Ñ  s   € ´ZÀÄ3Ó5G€ rJ   )Úparse_optionsr�  r9   r'  r  ÚKeyErrorr   r   r/   rÙ   Ú
parse_attr)rF   r$  rö   Ú
bound_argsr-  r  r  ÚsigkeysÚsigvalsr\   Úvr  Ú
constexprsÚpathÚattrvalsr  s                   rI   Ú
_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  rQ   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(  r…  r½  Ú_do_compileÚobjectrD   r  rV   rn   r,   Úlaunch_metadatarÙ   rî   ÚfunctionÚpacked_metadataÚlaunch_enter_hookÚlaunch_exit_hook)"rF   rñ   rò   r‹   rö   r�  r°  rž  Úkernel_cacher„  r¢   r$  r¬  r¶  r-  r  Úinspect_stages_keyÚinspect_stages_hashr  Úkernelr  rº  r  Únot_presentr9   r^   rq   Úglobals_dictÚnewValÚ	grid_sizeÚgrid_0Úgrid_1Úgrid_2rÉ  s"                                     rI   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àˆrJ   c                óT   — | j                   €| j                  S | j                  |«      S rL   )Ú_reprr=  )rF   r^   s     rI   r[  zJITFunction.repr  s"   € Ø $§
¡
Ð 2ˆt�}‰}ÐE¸¿
¹
À1»ÐErJ   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   )r7   r8   rg   r	  ÚversionrÎ   rÏ   rÚ  rÉ  Ú_triton_jit_function_registryr
  r±   r�  r‡  r  r&  rÙ   r9   r(  rÊ   r   r­  rÅ  rÑ  r¿  rR   Ú	arg_namesrÞ   rÌ   rº  r£  )rF   r
  rÜ  rÎ   rÏ   r¿  rR   r[  rÉ  r‰  rÐ   ÚdnsÚdns_oaÚprH   s                 €rI   r8   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   r9   zSpecialization data is for z but trying to preload for r  r  r¢   c                ó&  •— ‰j                   j                  | «      r‰j                  | «      S t        | t        «      rJd| v r‰j	                  t        | d   «      «      S d| v r%| d   }|t        v r	t        |   S t        d|› d�«      ‚t        | «      S )Nrd   rÿ   zUnable to resolve JITFunction z for preload)rÆ   Úis_dtyper4   r_  rd   rˆ  rÝ  rV   )r…   Újf_keyÚtls     €rI   Ú_decode_constantz-JITFunction.preload.<locals>._decode_constantI  s•   ø€ Ø�x‰x× Ñ  Ô'Ø—x‘x “Ð&Ü˜%¤Ô&Ø %Ñ'ØŸ<™<Ô(@ÀÀ{ÑASÓ(TÓUÐUØ! UÑ*Ø" >Ñ2�FØÔ!>Ñ>Ü<¸VÑDÐDÜ&Ð)GÈÀxÈ|Ð'\Ó]Ð]Ü+¨EÓ2Ð2rJ   r  r  r  r  r  T)rò   )r  Útriton.languager½   r   rª  rÃ  Úloadsr=  rV   ró   r‚  rÅ  r  r'  r_  r  rˆ  r4   r/   r³  rÇ  )rF   r˜  r  r�  Údeserialized_objr  r  r^   r¢   r$  Údeserialized_targetrç  r  r…   rº  r  r  r  r  r  ræ  s                       @rI   ÚpreloadzJITFunction.preload8  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ô	3ô 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  rF   rG   r¢   s   €€€€€rI   Úasync_compilez.JITFunction._do_compile.<locals>.async_compile{  s!   ø€ Ø—|‘| C°À×@PÑ@PÐ\d�|ÓeÐerJ   c                ót   •
— | ‰‰<   ‰j                  t        j                  j                  ‰‰‰	‰‰‰‰g‰
«	       y rL   )r¡  r   rÁ  Újit_post_compile_hook)rÑ  r  rº  r�  rÎ  r  r  rF   r  r¢   rò   s    €€€€€€€€€€rI   Úfinalize_compilez1JITFunction._do_compile.<locals>.finalize_compile~  s;   ø€ Ø$*�˜SÑ!Ø—‘¤§¡× CÑ CÀSÈ)ÐU[Ð]cÐeoØ '¨%¨°&õ:rJ   )r¢   r  )rÅ  r¡  r   rÁ  Újit_cache_hookr§  r   Úactive_moderw   r    r   Úsubmitrš  r  ró  )rF   r  r  r�  rº  r  r  rò   r^   r$  Ú
async_moderT   rñ  rô  rÑ  rð  rÎ  rG   r¢   s   ````````       @@@@rI   rÇ  zJITFunction._do_compilem  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ô-àˆrJ   c                ó   — t        d«      ‚)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rV   ©rF   r‹   rö   s      rI   Ú__call__zJITFunction.__call__Œ  s   € ÜÐWÓXÐXrJ   c                 ó   — y rL   rç   ©rF   ÚinstanceÚowners      rI   Ú__get__zJITFunction.__get__‘  s   € àrJ   c                 ó   — y rL   rç   rý  s      rI   r   zJITFunction.__get__•  s   € ð rJ   c                 ó   — y rL   rç   rý  s      rI   r   zJITFunction.__get__š  s   € ØrJ   c                óP   — d| j                   › d| j                  j                  › d�S )NzJITFunction(r   r  )r	  r
  r±   rN   s    rI   Ú__repr__zJITFunction.__repr__�  s&   € Ø˜dŸk™k˜]¨!¨D¯G©G×,@Ñ,@Ð+AÀÐCÐCrJ   )r¯   zbool | None)NNNNNNN)rF   z'JITFunction[Callable[P, R]]'r‹   úP.argsrö   úP.kwargsr¯   r%   rL   )rþ  r°   rÿ  úOptional[type]r¯   z'JITFunction[T]')rF   z-'JITFunction[Callable[Concatenate[U, P], R]]'rþ  r   rÿ  r  r¯   zCallable[P, R])rW   rg   r±   r�  r¡  r¤  r­  r½  rî   r[  r8   rì  rÇ  rû  r   r   r   r  r´   rµ   s   @rI   rþ   rþ   t  s�   ø„ òð.
ð 
ó.
ò`(ò/ò5ò07òrFð mqØ;?õ$ òL3
òjó>Yñ à	ó	ó 
ð	ð 
à,0ð	Ø)ð	Ø5Cò	ó 
ð	ó	öDrJ   rþ   c                 ó   — y rL   rç   r4  s    rI   Újitr	  ¦  s   € àrJ   ©rÜ  r[  rÉ  rÎ   rÏ   r¿  rR   c                 ó   — y rL   rç   r
  s          rI   r	  r	  «  s   € ð rJ   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¿  rR   r[  rÉ  )rn   r   rÁ  Ú	interpretÚinterpreterr  rþ   )	r
  r  r¿  rÎ   rÏ   rÉ  rR   r[  rÜ  s	     €€€€€€€rI   Ú	decoratorzjit.<locals>.decoratorÖ  sj   ø€ Ü˜Œ|Ðˆ|Ü�=‰=×"Ò"Ý8Ù& r°7ÐN_ØFdÐlqØ08¸tÐUdôfð fô ØØØ"3Ø/MØØ!ØØ /ô	ð 	rJ   ©r
  r#   r¯   zJITFunction[T]rç   )	r
  rÜ  r[  rÉ  rÎ   rÏ   r¿  rR   r  s	    ``````` rI   r	  r	  ¹  s&   þ€ ÷:ó ð& 
€~Ù˜‹}Ðð ÐrJ   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)rH   rW   rg   rô   )rŒ   s    rI   rõ   zMockTensor.wrap_dtypeû  s.   € à�=‰=×!Ñ! WÒ,°·±À7Ò1JÜ˜c“?Ð"Øˆ
rJ   Nc                ó*   — |€dg}|| _         || _        y )Nr   )rÆ   Úshape)rF   rÆ   r  s      rI   r8   zMockTensor.__init__  s   € Øˆ=Ø�CˆEØˆŒ
Øˆ�
rJ   c                ó†   — dg}| j                   dd  D ]  }|j                  |d   |z  «       Œ t        t        |«      «      S )Nr   rº   )r  r(  r‚  Úreversed)rF   ÚstridesÚsizes      rI   ÚstridezMockTensor.stride  sG   € Ø�#ˆØ—J‘J˜q˜r�Nò 	/ˆDØ�N‰N˜7 2™;¨Ñ-Õ.ð	/ä”X˜gÓ&Ó'Ð'rJ   c                  ó   — y©Nr   rç   rç   rJ   rI   Údata_ptrzMockTensor.data_ptr  ó   € àrJ   c                  ó   — yr  rç   rç   rJ   rI   Ú	ptr_rangezMockTensor.ptr_range  r   rJ   rL   )
rW   rg   r±   r²   Ústaticmethodrõ   r8   r  r  r"  rç   rJ   rI   rô   rô   õ  sM   „ ñð
 ñó ðó
ò(ð ñó ðð ñó ñrJ   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 rL   )rÆ   ÚbaseÚdatar�  r  )rF   r'  rÆ   s      rI   r8   zTensorWrapper.__init__  s5   € ØˆŒ
ØˆŒ	Ø—I‘IˆŒ	Ø—k‘kˆŒØ—Y‘Y—_‘_ˆ�
rJ   c                ó6   — | j                   j                  «       S rL   )r'  r  rN   s    rI   r  zTensorWrapper.data_ptr  s   € Ø�y‰y×!Ñ!Ó#Ð#rJ   c                ó4   —  | j                   j                  |Ž S rL   )r'  r  )rF   r‹   s     rI   r  zTensorWrapper.stride"  s   € Øˆt�y‰y×Ñ Ð&Ð&rJ   c                ó<   — d| j                   › d| j                  › d�S )NzTensorWrapper[r�  r  )rÆ   r'  rN   s    rI   Ú__str__zTensorWrapper.__str__%  s   € Ø §
¡
˜|¨2¨d¯i©i¨[¸Ð:Ð:rJ   c                ó6   — | j                   j                  «       S rL   )r'  Úelement_sizerN   s    rI   r.  zTensorWrapper.element_size(  s   € Ø�y‰y×%Ñ%Ó'Ð'rJ   c                ó^   — t        | j                  j                  «       | j                  «      S rL   )r%  r'  ÚcpurÆ   rN   s    rI   r0  zTensorWrapper.cpu+  s   € Ü˜TŸY™YŸ]™]›_¨d¯j©jÓ9Ð9rJ   c                óN   — | j                   j                  |j                   «       y rL   )r'  Úcopy_)rF   Úothers     rI   r2  zTensorWrapper.copy_.  s   € Ø�	‰	�‰˜Ÿ
™
Õ#rJ   c                ó^   — t        | j                  j                  «       | j                  «      S rL   )r%  r'  ÚclonerÆ   rN   s    rI   r5  zTensorWrapper.clone1  s   € Ü˜TŸY™YŸ_™_Ó.°·
±
Ó;Ð;rJ   c                ó`   — t        | j                  j                  |«      | j                  «      S rL   )r%  r'  ÚtorÆ   )rF   r�  s     rI   r7  zTensorWrapper.to4  s    € Ü˜TŸY™YŸ\™\¨&Ó1°4·:±:Ó>Ð>rJ   c                ó`   — t        | j                  j                  |«      | j                  «      S rL   )r%  r'  Ú	new_emptyrÆ   )rF   Úsizess     rI   r9  zTensorWrapper.new_empty7  s"   € Ü˜TŸY™Y×0Ñ0°Ó7¸¿¹ÓDÐDrJ   Nræ   )rW   rg   r±   r8   r  r  r,  r.  r0  r2  r5  r7  r9  rç   rJ   rI   r%  r%    s5   „ ò%ò$ò'ó;ò(ò:ò$ò<ò?óErJ   r%  c                óø   — t        | t        «      r;|| j                  j                  k(  r| j                  S t        | j                  |«      S t	        | d«      rt        | |«      S t        dt        | «      › d�«      ‚)Nr  zCannot reinterpret a r3  )r4   r%  r'  rÆ   Úhasattrr¦   rj   )ÚtensorrÆ   s     rI   Úreinterpretr>  ;  sk   € Ü�&œ-Ô(Ø�F—K‘K×%Ñ%Ò%à—;‘;Ðô ! §¡¨eÓ4Ð4Ü	�˜Ô	$ä˜V UÓ+Ð+äÐ/´°V³¨~¸QÐ?Ó@Ð@rJ   c                ó~   — |}t        | «      D ],  \  }}|j                  «       j                  d«      sŒ&||z  } |S  |S )Nzdef )r‡  r¿   rÀ   )r9  r:  rE  ÚidxÚlines        rI   rD  rD  J  sR   € Ø/Ðô ˜wÓ'ò ‰	ˆˆTØ�:‰:‹<×"Ñ" 6Õ*Ø  CÑ'Ð ØØÐð	ð  ÐrJ   c                óˆ   — t         j                  | «      }|st        d«      ‚t        |j	                  d«      «      }|dz  }|S )Nz'No function definition found for kernelÚindentr   )ÚINDENT_PATTERNrK  r<  r,   Úgroup)rR  Úindented_defrG  s      rI   rF  rF  X  sH   € ä!×(Ñ(¨Ó5€LÙÜÐBÓCÐCä˜l×0Ñ0°Ó:Ó;Ðà˜1ÑÐØÐrJ   c                  ó(   — e Zd Zd„ Zed„ «       Zd„ Zy)ÚBoundConstexprFunctionc                ó    — || _         || _        y rL   )Ú__self__Ú__func__)rF   rþ  r
  s      rI   r8   zBoundConstexprFunction.__init__f  s   € Ø ˆŒØˆ�rJ   c                ó.   — | j                   j                  S rL   )rK  rT   rN   s    rI   rT   z BoundConstexprFunction.cache_keyj  s   € à�}‰}×&Ñ&Ð&rJ   c                óB   —  | j                   | j                  g|¢­i |¤ŽS rL   )rK  rJ  rú  s      rI   rû  zBoundConstexprFunction.__call__n  s!   € Øˆt�}‰}˜TŸ]™]Ð<¨TÒ<°VÑ<Ð<rJ   N)rW   rg   r±   r8   r³   rT   rû  rç   rJ   rI   rH  rH  d  s    „ òð ñ'ó ð'ó=rJ   rH  c                  óB   ‡ — e Zd Zˆ fd„Zd„ Zedd„«       Zddœd„Zˆ xZS )ÚConstexprFunctionc                ó$   •— t         ‰| �  |«       y rL   )r7   r8   )rF   r
  rH   s     €rI   r8   zConstexprFunction.__init__t  s   ø€ Ü‰Ñ˜ÕrJ   c                ó"   — |�t        || «      S | S rL   )rH  )rF   r  Úobjclasss      rI   r   zConstexprFunction.__get__w  s   € àˆ?Ü)¨#¨tÓ4Ð4ØˆrJ   c                 ó   — y rL   rç   rú  s      rI   rû  zConstexprFunction.__call__}  s   € àrJ   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_constexprrd   )r¼   rV  rd   r  r
  r   rÁ  r  )
rF   rT  r‹   rö   rV  rd   r  r\   r¹  Úress
             rI   rû  zConstexprFunction.__call__�  sŽ   € ßHà15Ö6¨AÑ$ QÕ'Ð6ˆÐ6Ø;A¿<¹<»>×J±°!°Q�!Ñ)¨!Ó,Ñ,ÐJˆÑJð ˆd�g‰g�tÐ&˜vÑ&ˆàÐàˆJô �=‰=×"Ò"ØˆJÙ˜‹~Ðùò 7ùÛJs
   �B±B)rF   z#'ConstexprFunction[Callable[P, R]]'r‹   r  rö   r  r¯   r%   )rW   rg   r±   r8   r   r   rû  r´   rµ   s   @rI   rO  rO  r  s0   ø„ ôòð òó ðð )-÷ ð rJ   rO  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.
    )rO  r4  s    rI   Úconstexpr_functionrY  ”  s   € ô ˜RÓ Ð rJ   ræ   )Fr  )r[  úOptional[Callable]rÉ  rZ  rÎ   úOptional[Iterable[int | str]]rÏ   r[  r¿  úOptional[bool]rR   r\  r¯   zCallable[[T], JITFunction[T]]rL   )r
  zOptional[T]r[  rZ  rÉ  rZ  rÎ   r[  rÏ   r[  r¿  r\  rR   r\  r¯   zKernelInterface[T])r
  r#   r¯   zConstexprFunction[T])RÚ
__future__r   r   rz   r5   r:   rÔ   r”   r>  rJ  rH  Úcollectionsr   Údataclassesr   Ú	functoolsr   Útypingr   r	   r
   r   r   r   r   r   r   r   r   r   Útriton.backendsr   Útypesr   rh   r   r   r   Ú_utilsr   r   r   r   Úcacher   Útriton._C.libtritonr    r!   r"   rB   rA   rš  rL  rD  r#   r$   r%   r&   ÚNodeVisitorr(   rÂ   rÊ   rë   rí   r  r1  r5  rS   rÝ  r{  r…  rˆ  rþ   r	  rô   r%  r>  rD  rF  rH  rO  rY  rç   rJ   rI   ú<module>rh     sˆ  ðß ,Û 
Û Û Û Û Û Û 	Û Ý #Ý !Ý %÷ E÷  E÷  Eó  Eå 'Ý Ý Ý Ý ß `Ó `Ý  ß [Ñ [à!€Ø3€à�—‘Ð@À"Ç,Á,ÓO€áˆCƒL€ÙˆcƒN€ÙˆCƒL€ÙˆCƒL€ôf!˜Ÿ™ô f!ó\D÷4/>ñ />ódTôZ�g˜a‘jô Zò&ò$>*òB0÷r1ñ r1ðj !#Ð ð ÷ð ó ðòò0	ôjD�+˜¨qÑ1ô jDðd	 
òó 
ðð 
ð Ø#Ø*.Ø7;ØDHØ Ø#ñ
ð ð
ð (ð	
ð
 5ð
ð %Bð
ð ð
ð ð
ð #ò
ó 
ð
ð ð4ð Ø#Ø*.Ø7;ØDHØ Ø#ñ4Øð4ð ð	4ð
 (ð4ð 5ð4ð %Bð4ð ð4ð ð4ð ó4÷xñ ÷B"Eñ "EòJAò ò	ô=˜[ô =ô˜ W¨Q¡Zô ôD!rJ   