ó
    “~i|   ã            	      ó  • S SK JrJr  S SKrS SKrS SKrS SKrS SKrS SKrS SK	r	S SK
r
S SKJr  S SKJr  S SKJr  S SKJrJrJrJrJrJrJrJrJr  S SKJr  S SKJr  S	S
KJ r   SSK!J!r!  SSKJ"r"  S	SK#J$r$J%r%J&r&J'r'  SSK(J)r)  S SK*J+r+J,r,  Sr-Sr.\" S5      r/ " S S\R`                  5      r1S6S jr2 " S S5      r3S7S jr4 " S S\\/   5      r5S r6S r7S r8 " S S 5      r9\ " S! S"5      5       r:S# r;S$ r< " S% S&\9\5\/   5      r=\S8S' j5       r>\SSSSSSSS(.             S9S) jj5       r> S:SSSSSSSS(.               S;S* jjjr> " S+ S,5      r? " S- S.5      r@S/ rAS0 rB " S1 S2\95      rC " S3 S4\95      rDS5 rEg)<é    )ÚannotationsÚdivisionN)Údefaultdict)Ú	dataclass)Úcached_property)	ÚCallableÚGenericÚIterableÚOptionalÚTypeVarÚoverloadÚDictÚAnyÚTuple)ÚBaseBackend)Ú
ModuleTypeé   )Úknobsé   )Údriver)Ú_async_compile)Úfind_paths_ifÚget_iterable_pathÚtype_canonicalisation_dictÚis_namedtuple)Úget_cache_key)Úget_cache_invalidating_env_varsÚnative_specialize_implztriton.languagez"triton.experimental.gluon.languageÚTc                  óŒ   ^ • \ rS rSrSrSU 4S jjr\S 5       rS rS r	SS jr
S rS	 rS
 rS rS rS rS rS rS rSrU =r$ )ÚDependenciesFinderé"   a±  
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U ]  5         Xl        [        R                  " UR                  S5      5      U l        X l        X0l        1 SkU l	        [        [        SS1U l        0 U l        SU l        g )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        €ÚO/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/runtime/jit.pyr2   ÚDependenciesFinder.__init__.   sn   ø€ Ü‰ÑÔØŒ	Ü—n’n S§Z¡Z°Ó%8Ó9ˆŒð ŒØ"Œò*
ˆÔ&ô ÜØØð	"
ˆÔð" TVˆÔà*/ˆÕ'ó    c                ó6   • U R                   R                  5       $ ©N)r7   Ú	hexdigest©r@   s    rC   ÚretÚDependenciesFinder.retY   s   € à�{‰{×$Ñ$Ó&Ð&rE   c                ó’   • [         R                  " UR                  5      (       a  g[        USS5      nUR	                  [
        5      $ )NTÚ
__module__Ú )ÚinspectÚ	isbuiltinÚfuncr-   Ú
startswithr<   )r@   ÚnoderQ   Úmodules       rC   Ú_is_triton_builtinÚ%DependenciesFinder._is_triton_builtin]   s9   € Ü×Ò˜TŸY™Y×'Ñ'ØÜ˜˜|¨RÓ0ˆØ× Ñ ¤Ó/Ð/rE   c                óF  • [        U[        5      (       d   eU R                  R                  5       UR                  R                  5       -   H]  nUu  p4U R                  U   u  pTUR                  U   u  pdXV:w  d  M0  [	        SU SU SU R
                   SUR                   SU S35      e   U R                  R                  UR                  5        UR                  nU[        [        USS5      5      -  nU R                  R                  UR                  S	5      5        g )
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@   rQ   ÚkÚvar_nameÚ_Úv1Úv2Úfunc_keys           rC   Ú_update_hashÚDependenciesFinder._update_hashc   s&  € Ü˜$¤×,Ñ,Ð,Ð,ð ×&Ñ&×+Ñ+Ó-°×0EÑ0E×0JÑ0JÓ0LÔLˆAØ‰KˆHØ×)Ñ)¨!Ñ,‰EˆBØ×)Ñ)¨!Ñ,‰EˆBØ�xÜ"Ø& x j°¸B¸4Ð?OÐPT×PYÑPYÈ{ÐZmÐnr×n{Ñn{Ðm|ð  }Tð  UWð  TXð  XOð  Póð ñ Mð 	×Ñ×$Ñ$ T×%:Ñ%:Ô;à—>‘>ˆØ”Cœ  j°%Ó8Ó9Ñ9ˆØ�‰×Ñ˜8Ÿ?™?¨7Ó3Õ4rE   c                óR  • SSK Jn  Ub  [        U5      [        L a  g [	        USS5      (       a%  UR
                   H  nU R                  U5        M     g [	        USS5      (       a  g [	        USS5      S:X  a  g [        U[        5      (       a  U R                  U5        g [        U5      (       a3  [        U[        5      (       d  [        X5      (       d  [        S	U 35      eU R                  (       a  g Ub0  [        R                  " U5      U4U R                  U[!        U5      4'   g )
Nr   ©Ú	constexprÚ__triton_aggregate__FÚ__triton_builtin__rM   rN   ztriton.language.extra.libdevicez!Unsupported function referenced: )Úlanguage.corerk   Útyper   r-   Ú
hash_attrsÚrecord_referencer.   rZ   rg   Úcallabler\   r?   r/   Údeepcopyr>   Úid)r@   ÚvalÚvar_dictr3   rk   Úattrs         rC   rq   Ú#DependenciesFinder.record_referenceu   s  € Ý-ð ‰;œ$˜s›)¤zÒ1Øä�3Ð.°×6Ñ6ØŸœ�Ø×%Ñ% dÖ+ñ 'àä�3Ð,¨e×4Ñ4Øô �3˜ bÓ)Ð-NÓNØä�cœ;×'Ñ'Ø×Ñ˜cÔ"Øä�C�=‰=¤¨C´×!6Ñ!6¼zÈ#×?YÑ?YÜÐ!BÀ3À%ÐHÓIÐIð ×*×*ØàÑÜ;?¿=º=ÈÓ;MÈxÐ:XˆD×!Ñ! 4¬¨H«Ð"6Ñ7ØrE   c                óF  ^ • [        UR                  5      [        R                  L a  UR                  $ UR                  T R
                  ;   a  g U 4S jnU" UR                  5      u  p4UR                  T R                  ;   a  U$ T R                  X4UR                  5        U$ )Nc                óº   >• TR                   R                  U S 5      nUb  UTR                   4$ TR                  R                  U S 5      nUb  UTR                  4$ g)N©NN)r8   Úgetr9   )r3   ru   r@   s     €rC   Úname_lookupÚ2DependenciesFinder.visit_Name.<locals>.name_lookup¤   sZ   ø€ Ø—,‘,×"Ñ" 4¨Ó.ˆCØ‰Ø˜DŸL™LÐ(Ð(Ø—.‘.×$Ñ$ T¨4Ó0ˆCØ‰Ø˜DŸN™NÐ*Ð*ØrE   )ro   ÚctxÚastÚStorert   Úlocal_namesr:   rq   )r@   rS   r}   ru   rv   s   `    rC   Ú
visit_NameÚDependenciesFinder.visit_Nameœ   s|   ø€ Ü�—‘‹>œSŸY™YÒ&Ø—7‘7ˆNà�7‰7�d×&Ñ&Ó&àõ	ñ $ D§G¡GÓ,‰ˆØ�7‰7�d×4Ñ4Ó4ØˆJà×Ñ˜c¨T¯W©WÔ5Øˆ
rE   c                ób   • UR                    Vs/ s H  o R                  U5      PM     sn$ s  snf rG   )ÚeltsÚvisit)r@   rS   Úelts      rC   Úvisit_TupleÚDependenciesFinder.visit_Tuple´   s&   € ð ,0¯9ª9Ó5ª9 C—
‘
˜3–©9Ñ5Ð5ùÒ5s   �,c                ó~  • U R                  UR                  5      n[        U[        R                  5      (       a<  U R                  UR                  5      n[        U[        R                  5      (       a  M<  [        USS5      nUb  X0R                  ;   a  g [        X!R                  5      nU R                  U5        U$ )Nr]   rN   )	r‡   Úvaluer.   r€   Ú	Attributer-   r=   rw   rq   )r@   rS   ÚlhsÚlhs_namerJ   s        rC   Úvisit_AttributeÚ"DependenciesFinder.visit_Attribute¹   sŒ   € Ø�j‰j˜Ÿ™Ó$ˆÜ˜œcŸm™m×,Ñ,Ø—*‘*˜SŸY™YÓ'ˆCô ˜œcŸm™m×,Ó,ä˜3 
¨BÓ/ˆØ‰;˜(×&<Ñ&<Ó<ØÜ�cŸ9™9Ó%ˆØ×Ñ˜cÔ"Øˆ
rE   c                óš   • UR                   R                    Vs1 s H  o"R                  iM     snU l        U R                  U5        g s  snf rG   )ÚargsÚargr‚   Úgeneric_visit)r@   rS   r”   s      rC   Úvisit_FunctionDefÚ$DependenciesFinder.visit_FunctionDefÄ   s6   € à/3¯y©y¯~ª~Ó>ª~¨ŸGœG©~Ñ>ˆÔØ×Ñ˜4Õ ùò ?s   ™Ac                ó„  ^ • U 4S jn[         R                  " UR                  UR                  UR                  (       a  UR                  /O/ UR
                  5       H  nT R                  U5        M     U" UR                  5        UR                  b  T R                  UR                  5        U" UR                  5        g )Nc                ó    >•  TR                   (       a   eSTl         U  H  nUc  M  TR                  U5        M     STl         g ! STl         f = f)NTF)r?   r‡   )ÚdefaultsÚexprr@   s     €rC   Úvisit_defaultsÚ:DependenciesFinder.visit_arguments.<locals>.visit_defaultsÕ   sN   ø€ ð8Ø×:×:Ð:Ð:Ø26�Ô/Û$�DØÓ'ØŸ
™
 4Ö(ñ %ð 38�Õ/ø°%�Õ/ús   ƒ A §A Á	A)
Ú	itertoolsÚchainÚposonlyargsr“   ÚvarargÚ
kwonlyargsr‡   Úkw_defaultsÚkwargrš   )r@   rS   rœ   r”   s   `   rC   Úvisit_argumentsÚ"DependenciesFinder.visit_argumentsÉ   sˆ   ø€ õ	8ô —?’? 4×#3Ñ#3°T·Y±YÐQU×Q\×Q\ÀÇÁÁÐbdÐfj×fuÑfuÖvˆCØ�J‰J�sŽOñ wñ 	�t×'Ñ'Ô(à�:‰:Ñ!Ø�J‰J�t—z‘zÔ"á�t—}‘}Õ%rE   c                óÄ   • U R                  U5      n[        U[        5      (       a  U =R                  [	        U5      -  sl        g U R                  R                  U5        g rG   )r‡   r.   r)   r‚   ÚsetÚadd)r@   rS   Útargets      rC   ÚvisitAssnTargetÚ"DependenciesFinder.visitAssnTargeté   sH   € ð —‘˜DÓ!ˆÜ�fœd×#Ñ#Ø×Ò¤ F£Ñ+Öà×Ñ× Ñ  Õ(rE   c                óª   • [        UR                  5      S:w  a  [        S5      eU R                  UR                  S   5        U R	                  U5        g )Nr   z2Simultaneous multiple assignment is not supported.r   )r&   ÚtargetsÚ	TypeErrorr«   r•   ©r@   rS   s     rC   Úvisit_AssignÚDependenciesFinder.visit_Assignò   sG   € Üˆt�|‰|Ó Ó!ô
 ÐPÓQÐQà×Ñ˜TŸ\™\¨!™_Ô-ð 	×Ñ˜4Õ rE   c                ó\   • U R                  UR                  5        U R                  U5        g rG   ©r«   rª   r•   r°   s     rC   Úvisit_AnnAssignÚ"DependenciesFinder.visit_AnnAssignÿ   ó$   € Ø×Ñ˜TŸ[™[Ô)ð 	×Ñ˜4Õ rE   c                ó\   • U R                  UR                  5        U R                  U5        g rG   r´   r°   s     rC   Ú	visit_ForÚDependenciesFinder.visit_For  r·   rE   )	r8   r7   r‚   r3   r9   r=   r:   r>   r?   )ÚreturnÚNoner{   )r]   rM   Ú__qualname__Ú__firstlineno__Ú__doc__r2   ÚpropertyrJ   rU   rg   rq   rƒ   r‰   r�   r–   r¥   r«   r±   rµ   r¹   Ú__static_attributes__Ú__classcell__©rB   s   @rC   r!   r!   "   se   ø† ñ	÷)0ðV ñ'ó ð'ò0ò5ô$%òNò06ò
	ò!ò
&ò@)ò!ò!÷!ð !rE   r!   c                óZ  • SS K Js  Jn  [        U [        5      (       aá  U R                  5       n U R                  S5      (       a<  U R                  S5      n [        U 5      n U R                  S5      (       d   eSU SS  -   $ U R                  S5      (       a  S[        U S S 5      -   $ U R                  S5      (       a  S[        U SS  5      -   $ U R                  S5      (       a  [        U R                  S5      5      $ O†[        XR                  5      (       a  S[        U R                  5       3$ [        XR                  5      (       a  U R                  n O-[        U [        5      (       a  U R                  n O[	        U 5      n [         R"                  " U R%                  S	S
5      U 5      $ )Nr   zconst ÚconstÚ*ú*kr   éÿÿÿÿztl.Ú_trN   )Útriton.language.coreÚlanguageÚcorer.   r`   ÚstriprR   ÚremoveprefixÚ_normalize_tyÚendswithÚpointer_typeÚ
element_tyÚdtyper3   ro   r]   r   r|   Ú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ÐCrE   c                  óž   • \ rS rSrSr  SS jr\S 5       r\SS j5       r\SS j5       r	\S 5       r
\S 5       r\S	 5       r\S
 5       rSrg)ÚKernelParami+  zBRepresents a parameter (name plus metadata) to a @jit'ed function.c                ó4   • Xl         X l        X0l        X@l        g rG   )ÚnumÚ_paramÚdo_not_specializeÚdo_not_specialize_on_alignment)r@   rÙ   ÚparamrÛ   rÜ   s        rC   r2   ÚKernelParam.__init__.  s   € àŒØŒØ!2ÔØ.LÕ+rE   c                ó.   • U R                   R                  $ rG   )rÚ   r3   rI   s    rC   r3   ÚKernelParam.name5  s   € à�{‰{×ÑÐrE   c                óÜ   • U R                   R                  (       a2  U R                   R                  [        R                  R                  :X  a  g[        U R                   R                  5      $ )NrN   )rÚ   Ú
annotationrO   Ú	ParameterÚemptyrÏ   rI   s    rC   râ   ÚKernelParam.annotation9  sD   € à�{‰{×%×%¨¯©×)?Ñ)?Ä7×CTÑCT×CZÑCZÓ)ZØÜ˜TŸ[™[×3Ñ3Ó4Ð4rE   c                óæ   • U R                   nUR                  S5      (       a  USS  nOUR                  S5      (       a  USS  nU[        [        R                  " 5       5      ;   a  U R                   $ g)NrÇ   r   rÆ   r   rN   )râ   rR   r¨   r   Úvalues)r@   Úas     rC   Úannotation_typeÚKernelParam.annotation_type?  sc   € à�O‰OˆØ�<‰<˜×ÑØ�!�"�‰AØ�\‰\˜#×ÑØ�!�"�ˆAØ”Ô.×5Ò5Ó7Ó8Ó8Ø—?‘?Ð"ØrE   c                ó    • SU R                   ;   $ ©Nrk   )râ   rI   s    rC   Úis_constexprÚKernelParam.is_constexprJ  s   € à˜dŸo™oÑ-Ð-rE   c                óˆ   • U R                   (       a  gSU R                  ;   =(       d    U R                  R                  S5      $ )NFrÅ   rÇ   )rí   râ   rR   rI   s    rC   Úis_constÚKernelParam.is_constN  s1   € à××ØØ˜$Ÿ/™/Ñ)×M¨T¯_©_×-GÑ-GÈÓ-MÐMrE   c                ó.   • U R                   R                  $ rG   )rÚ   ÚdefaultrI   s    rC   ró   ÚKernelParam.defaultT  s   € à�{‰{×"Ñ"Ð"rE   c                ód   • U R                   R                  [        R                  R                  :g  $ rG   )rÚ   ró   rO   rã   rä   rI   s    rC   Úhas_defaultÚKernelParam.has_defaultX  s#   € à�{‰{×"Ñ"¤g×&7Ñ&7×&=Ñ&=Ñ=Ð=rE   )rÚ   rÛ   rÜ   rÙ   N)rÙ   r%   rÝ   zinspect.ParameterrÛ   ÚboolrÜ   rø   ©r»   r`   )r]   rM   r½   r¾   r¿   r2   r   r3   râ   ré   rí   rð   rÀ   ró   rö   rÁ   © rE   rC   r×   r×   +  s¢   † ÙLðMØ15ôMð ñ ó ð ð ó5ó ð5ð
 óó ðð ñ.ó ð.ð ñNó ðNð
 ñ#ó ð#ð ñ>ó ó>rE   r×   c                ó2   • SnSn[        [        XX5      S   $ )NFTr   )r   r   )r”   Ú
specializerð   Úaligns       rC   Úmangle_typerþ   ]  s!   € Ø€HØ€EÜ!¤+¨s¸jÓPÐQRÑSÐSrE   c                  ó6   • \ rS rSr% S\S'   S rS rS	S jrSrg)
ÚKernelInterfaceic  r   Úrunc               ó\   • U R                   " [        [        R                  U5      USS.UD6$ )NT©ÚgridÚwarmup)r  ÚmapÚ
MockTensorÚ
wrap_dtype)r@   r  r“   Úkwargss       rC   r  ÚKernelInterface.warmupf  s(   € Ø�xŠx´´Z×5JÑ5JÈDÓ1Q˜T¨$Ñ\ÐU[Ñ\Ð\rE   c               ó   • [        S5      e)Nzrun not implemented)ÚNotImplementedError)r@   r  r  r“   r	  s        rC   r  ÚKernelInterface.runi  s   € Ü!Ð"7Ó8Ð8rE   c                ó   ^ ^• UU 4S j$ )z�
A JIT function is launched with: fn[grid](*args, **kwargs).
Hence JITFunction.__getitem__ returns a callable proxy that
memorizes the grid.
c                 ó.   >• TR                   " U TSS.UD6$ )NFr  )r  )r“   r	  r  r@   s     €€rC   Ú<lambda>Ú-KernelInterface.__getitem__.<locals>.<lambda>r  s   ø€  t§x¢xÈ$°TÀ%Ñ'YÐRXÒ'YrE   rú   )r@   r  s   ``rC   Ú__getitem__ÚKernelInterface.__getitem__l  s   ù€ õ ZÐYrE   rú   N)r»   r   )	r]   rM   r½   r¾   Ú__annotations__r  r  r  rÁ   rú   rE   rC   r   r   c  s   ‡ Ø	
ƒFò]ò9÷ZrE   r   c           
     óV  • UR                  5        VVs0 s HT  u  pVXVR                  R                  S:X  a  [        U5      O)UR                  R                  S:X  a  SUR                  0OU_MV     nnnSS KnXUR                  5        Vs/ s H  n[        U5      PM     sn[        UR                  5       5      UR                  5        Vs/ s H  n[        U5      PM     sn[        UR                  5       5      UR                  WS.n	UR                  U	5      n
U
$ s  snnf s  snf s  snf )NrÓ   rk   r   )r3   Ú	signatureÚconstant_keysÚconstant_valsÚ
attrs_keysÚ
attrs_valsÚoptionsÚkey)ÚitemsrB   r]   r`   rŒ   Újsonr[   r)   rç   Ú__dict__Údumps)r3   r  Ú	constantsÚattrsr  r  rŒ   r  ÚxÚobjÚserialized_objs              rC   Úserialize_specialization_datar&  v  s  € ð $Ÿ/™/Ô+ôò ,‰JˆCð 	Ÿ?™?×3Ñ3°wÓ>ŒS�ŒZØ&+§o¡o×&>Ñ&>À+Ó&Mˆ�e—k‘kÑ"ÐSXò	Yá+ð ñ ó àÐQZ×Q_ÑQ_ÔQaÓ?bÒQaÈAÄÀQÆÑQaÑ?bÜˆY×ÑÓÓ À%Ç*Á*Ä,Ó0OÂ,¸Q´°a¶Á,Ñ0OÔ_cÐdi×dpÑdpÓdrÓ_sØ×#Ñ#¨Cñ€Cð
 —Z‘Z “_€NØÐùóùò @cùÚ0Os   ”ADÂ	D!Ã
D&c                ó`  • [        U R                  5      [        U5      :X  d   e/ n[        U R                  R                  5       U5       GH0  u  pEUR                  (       a  UR                  SU S35        M.  UR                  (       a  SOSnUR                  (       a  SOSnUR                  (       a  SOSnSU SU SU SU S3	n	UR                  (       a‘  [        UR                  [        5      (       a%  UR                  S:X  d  UR                  SS	 S
;   a  SnU(       a$  UR                  SUR                   SU	 S35        Mü  UR                  SUR                   S35        GM  UR                  U	 5        GM3     S n
SSR                  [        [        X R                  R                  5       5      5      S/-   5       SSR                  U R                  R                  5        Vs/ s H  nSU SU 3PM     sn5       SSR                  U5       S3nU R                  R                  5        VVs0 s H>  u  pLUR                   ["        R$                  R&                  Ld  M.  SU 3UR                   _M@     nnn[(        nXíS'   X-S'   [*        US'   [-        X½5        US   $ s  snf s  snnf )a  
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                óz   • U S   R                   [        R                  R                  L a  U S   $ U S    SU S    3$ )Nr   r   z	=default_)ró   rO   rã   rä   )r#  s    rC   r  Ú0create_function_from_signature.<locals>.<lambda>¨  sC   € ˜A˜a™DŸL™L¬G×,=Ñ,=×,CÑ,CÒC�A�a‘DÐaÈAÈaÉDÈ6ÐQZÐ[\Ð]^Ñ[_ÐZ`ÐIaÐarE   z
def dynamic_func(z	**optionsz):
    params = {Ú'z': z}
    specialization = [Ú,z-]
    return params, specialization, options
Údefault_Úspecialize_implÚbackendrZ   Údynamic_func)r&   Ú
parametersÚzipr[   rí   Úappendrð   rÛ   rÜ   ré   r.   r`   Újoinr)   r  r  ró   rO   rã   rä   r   rZ   Úexec)ÚsigÚkparamsr5  Úspecializationr3   Úkprð   rü   rý   rJ   r”   Ú	func_bodyrÝ   Úfunc_namespacer4  s                  rC   Úcreate_function_from_signaturerB  ‡  sƒ  € ô ˆs�~‰~Ó¤# g£,Ó.Ð.Ð.à€Nä˜Ÿ™×+Ñ+Ó-¨w×7‰ˆØ�?�?Ø×!Ñ! 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Ñ6Ø×)Ñ)¨TÓ1°R×5GÑ5GÈÈÐ5KÈ|Ó5[à%*˜
ÞØ"×)Ñ)¨B¨r×/AÑ/AÐ.BÀ&ÈÈÈTÐ*RÖSð #×)Ñ)¨B¨r×/AÑ/AÐ.BÀ(Ð*K×Là×%Ñ%¨¨×/ñ' 8ñ, b€CðØ—)‘)œD¤ S¯.©.×*>Ñ*>Ó*@Ó!AÓBÀkÀ]ÑRÓSÐTð UØ—	‘	¸3¿>¹>×;NÑ;NÔ;PÓQÒ;P°4˜Q˜t˜f C¨ vÓ.Ñ;PÑQÓRÐSð TØ—x‘x Ó/Ð0ð 1ð€Ið Ÿ>™>×/Ñ/Ô1ôâ1‰KˆDØ�=‰=¤× 1Ñ 1× 7Ñ 7Ð7ó 	)ˆ(�4�&Ð˜5Ÿ=™=Ò(Ù1ð ñ ô -€OØ(7Ð$Ñ%Ø '�9ÑÜ$/€N�=Ñ!ô 	ˆÔ#ð ˜.Ñ)Ð)ùò+ Rùós   Ç+J%È6-J*É'J*c                ó8   • U R                    SU R                   3$ )NÚ.)rM   r½   ©Úfns    rC   Úget_full_namerG  Ã  s   € Ø�m‰mˆ_˜A˜bŸo™oÐ.Ð/Ð/rE   c                  óp   • \ rS rSrS rS r\SS j5       rS rS r	\S 5       r
S rS	 rS
 r\" \\S9rSrg)rZ   iÇ  c                ó¼  • Xl         [        R                  " U5      U l         [        R                  " U5      u  U l        U l        [        U5      U l	        [        R                  " 5       U l        [        R                  " SR                  U R                  5      5      nU[         R"                  " SU[         R$                  5      R'                  5       S  nX0l        S U l        0 U l        UR.                  U l        UR0                  U l        UR2                  U l        UR4                  U l        UR6                  U l        g ! [         a  n[        S5      UeS nAff = f)Nz1@jit functions should be defined in a Python filerN   z^def\s+\w+\s*\()rF  rO   r  ÚgetsourcelinesÚraw_srcÚstarting_line_numberÚOSErrorÚ
ValueErrorrG  Ú_fn_nameÚ	threadingÚRLockÚ
_hash_lockÚtextwrapÚdedentr:  ÚreÚsearchÚ	MULTILINEÚstartÚ_srcÚhashr>   r¿   r]   r½   Ú__globals__rM   )r@   rF  ÚerA   s       rC   r2   Ú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   £#E  Å 
EÅ
EÅEc                óp   • U R                   [        R                  " U R                  5      R                  -  $ rG   )r[  rO   ÚgetclosurevarsrF  r9   rI   s    rC   Úget_capture_scopeÚJITCallable.get_capture_scopeë  s(   € Ø×Ñ¤'×"8Ò"8¸¿¹Ó"A×"KÑ"KÑKÐKrE   c                óæ  • U R                      U R                  b  U R                  sS S S 5        $ SU R                   3U l        [        R                  " U R
                  5      R                  n[        U R                  U R                  UU R                  S9nUR                  U R                  5       5        UR                  [        U R                  5      -   U l        [        [!        UR"                  R%                  5       5      5      U l        SSKJn  U =R                  [        U R"                  R%                  5        VVVs/ s H  u  u  pEu  pe[+        Xc5      (       d  M  XF4PM!     snnn5      -  sl        [,        R.                  " U R                  R1                  S5      5      R3                  5       U l        S S S 5        U R                  $ s  snnnf ! , (       d  f       U R                  $ = f)Nz
recursion:)r3   r8   r9   rA   r   rj   r$   )rR  rZ  rO  rO   r_  rF  r9   r!   r[  rA   r‡   ÚparserJ   r`   rL  ÚdictÚsortedr>   r  rÊ   rk   r.   r4   r5   r6   rH   )r@   r9   Údependencies_finderrk   r3   rc   ru   s          rC   r_   ÚJITCallable.cache_keyî  st  € ð �_‹_Ø�y‰yÑ$Ø—y‘y÷ ‰_ð
 % 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õ=Ú9VÑ"5¡) 4©X¨cÜ!+¨C×!;ó *˜t›kÙ9Vó=ó >ñ >�Iô  Ÿš t§y¡y×'7Ñ'7¸Ó'@ÓA×KÑKÓMˆDŒI÷# ð$ �y‰yÐùô	=÷ Œ_ð$ �y‰yÐús*   �G°DGÅGÅ#GÅ*AGÇGÇ
G0c                ó,   • [        U R                  5      $ rG   )rZ  r_   rI   s    rC   Ú__hash__ÚJITCallable.__hash__  s   € Ü�D—N‘NÓ#Ð#rE   c                ó  • [         R                  " U R                  5      n[        U[         R                  5      (       d   e[        UR                  5      S:X  d   e[        UR                  S   [         R                  5      (       d   eU$ )Nr   r   )r€   rc  rY  r.   ÚModuler&   ÚbodyÚFunctionDef)r@   Útrees     rC   rc  ÚJITCallable.parse  se   € Ü�yŠy˜Ÿ™Ó#ˆÜ˜$¤§
¡
×+Ñ+Ð+Ð+Ü�4—9‘9‹~ Ó"Ð"Ð"Ü˜$Ÿ)™) A™,¬¯©×8Ñ8Ð8Ð8ØˆrE   c                ó   • SSK Jn  U" U 5      $ )Nr   )Úconstexpr_type)rÊ   rr  )r@   rr  s     rC   ro   ÚJITCallable.type  s   € å7Ù˜dÓ#Ð#rE   c                ó   • SU l         Xl        g)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)rZ  rY  )r@   Únew_srcs     rC   Ú_unsafe_update_srcÚJITCallable._unsafe_update_src  s   € ð ˆŒ	Ø�	rE   c                ó   • [        S5      e)NzqCannot set attribute 'src' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)ÚAttributeErrorrI   s    rC   Ú_set_srcÚJITCallable._set_src!  s   € Üð (ó )ð 	)rE   c                ó   • U R                   $ rG   )rY  rI   s    rC   Ú_get_srcÚJITCallable._get_src&  s   € Ø�y‰yÐrE   )ÚfgetÚfset)r¿   r[  rM   r]   r½   rO  rR  rY  rF  rZ  rK  r  rL  r>   Nrù   )r]   rM   r½   r¾   r2   r`  rÀ   r_   ri  rc  ro   rv  rz  r}  rA   rÁ   rú   rE   rC   rZ   rZ   Ç  s]   † ò (òDLð óó ðò,$òð ñ$ó ð$òò)ò
ñ ˜ xÑ
0ƒCrE   rZ   c                  ó4   • \ rS rSr% S\S'   S\S'   S\S'   Srg	)
ÚJitFunctionInfoi,  r   rT   r`   r3   ÚJITFunctionÚjit_functionrú   N)r]   rM   r½   r¾   r  rÁ   rú   rE   rC   r‚  r‚  ,  s   ‡ àÓØ
ƒIØÖrE   r‚  c                ó°   ^• [        U5      [        U5      4nU R                  US 5      nUb  U$ U4S jm[        T" U5      5      [        U5      -   nX@U'   U$ )Nc                ó|  >• [        U [        5      (       a  U  Vs/ s H  nT" U5      PM     sn$ [        U 5      (       a'  U  Vs/ s H  nT" U5      PM     nnU R                  " U6 $ [        U [        5      (       a  [	        U4S jU  5       5      $ [        U [
        5      (       a  U R                  $ U $ s  snf s  snf )Nc              3  ó4   >#   • U  H  nT" U5      v •  M     g 7frG   rú   )Ú.0r”   Úreplace_callabless     €rC   Ú	<genexpr>Ú?compute_cache_key.<locals>.replace_callables.<locals>.<genexpr>A  s   øé € Ð?º3°CÑ*¨3×/Ð/º3ùs   ƒ)r.   r)   r   rB   ÚtuplerZ   r_   )r$  r”   Úresultsr‰  s      €rC   r‰  Ú,compute_cache_key.<locals>.replace_callables:  s¡   ø€ Ü�cœ4× Ñ Ù69Ó:²c¨sÑ% cÖ*±cÑ:Ð:Ü˜3×ÑÙ9<Ó=º°#Ñ(¨Ö-¹ˆGÐ=Ø—=’= 'Ð*Ð*Ü˜œU×#Ñ#ÜÔ?¹3Ó?Ó?Ð?Ü˜œ[×)Ñ)Ø—=‘=Ð Øˆ
ùò ;ùâ=s   ›B4ÁB9)rŒ  r`   r|   )Úkernel_key_cacher>  r  r  r_   r‰  s        @rC   Úcompute_cache_keyr�  3  sb   ø€ Ü�Ó ¤# g£,Ð
/€CØ ×$Ñ$ S¨$Ó/€IØÑØÐõ
ô Ñ% nÓ5Ó6¼¸W»ÑE€IØ%�SÑØÐrE   c                óˆ   • [        U [        5      (       d  U $ [        U 5       H  u  p[        U5      X'   M     [	        U 5      $ rG   )r.   r)   Ú	enumerateÚconvert_to_tuple_if_listrŒ  )ÚitemÚiÚnested_values      rC   r“  r“  K  s>   € ä�dœD×!Ñ!Øˆô % Tž?‰ˆÜ*¨<Ó8ˆ‹ñ +ô �‹;ÐrE   c                  óz   ^ • \ rS rSrS r  SS jrS rS rS rS r	S r
  SU 4S	 jjrS
 rS rS rS rSrU =r$ )rƒ  iW  c                ó   • g)NFrú   rI   s    rC   Úis_gluonÚJITFunction.is_gluonY  s   € ØrE   c	                ó²  • U(       d  g U R                   R                  n	U R                   R                  n
SR                  [	        U R
                  US   5       VVs/ s H  u  p¼UR                   SU 3PM     snn5      nU	 SUR                   SUR                   SUR                   SUR                   SUR                   S	U S
3n[        U R                   5      n[        XóXWS   Xb5      nUUUUR                  UR                  UR                  UR                  UR                  UR                  UUUS.nU" UU[        X©U 5      SU0UEUSS9$ s  snnf )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  ÚreprrF  ÚcompileÚis_manual_warmupÚalready_compiled)rF  r½   rM   r:  r8  Úparamsr3   rž  rŸ  r   r¡  r¢  rG  r&  r£  r‚  )r@   Úhookr  r  r�  r!  r  r¤  r¦  r3   rT   rÝ   rÕ   Ú	arg_reprsr§  Ú	full_namer¥  r	  s                     rC   Ú
_call_hookÚJITFunction._call_hook\  s°  € ö Øà�w‰w×#Ñ#ˆØ—‘×#Ñ#ˆØ—I‘IÄcÈ$Ï+É+ÐWZÐ[\ÑW]ÔF^Ô_ÒF^¹¸ %§*¡* ¨R°¨tÓ4ÑF^Ò_Ó`ˆ	Ø��{ 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ÐR[ÐefÑ]gÐipÓvÐð #ØØ"Ø ×*Ñ*Ø×(Ñ(Ø!×,Ñ,Ø '× 8Ñ 8Ø'.×'FÑ'FØ"×.Ñ.ØØ#6Ø"ñ
ˆñ ØØÜ˜v¨TÓ2Ø˜CÐ* 6Ð*Ø&Ø"ñ
ð 	
ùó+ `s   ÁE
c                ó^   • [        U5      (       d   eU R                  R                  U5        g)zu
Add a hook that will be executed prior to the execution of run
function with args and kwargs passed into the kernel
N)rr   Úpre_run_hooksr9  )r@   r¬  s     rC   Úadd_pre_run_hookÚJITFunction.add_pre_run_hookŠ  s&   € ô
 ˜�~‰~Ðˆ~Ø×Ñ×!Ñ! $Õ'rE   c                óØ   • SSK JnJnJnJn  [
        R                  R                  5       nU" U5      nXl        X l        X0l        [        U R                  U R                  U5      n0 0 XVU4$ )z!
Precompute as much as possible.
r   )ÚCompiledKernelr¨  Ú	ASTSourceÚmake_backend)Úcompilerr¶  r¨  r·  r¸  r   ÚactiveÚget_current_targetrB  r  r«  )r@   r¶  r¨  r·  r¸  rª   r5  Úbinders           rC   Úcreate_binderÚJITFunction.create_binder’  s\   € ÷ 	PÓOÜ—‘×1Ñ1Ó3ˆÙ˜vÓ&ˆØ,ÔØŒØ"ŒÜ/°·±ÀÇÁÈWÓUˆØ�2�v¨Ð.Ð.rE   c           
     óØ  • UR                  U5      nU R                   Vs/ s H  ofR                  PM     nnU Vs/ s H  ofS   PM	     nn[        Xx5       V	V
s0 s H  u  pšXš_M	     nn	n
SU;  d   S5       eSU;  d   S5       eSU;  d   S5       eU H'  n	X•R                  ;  d  M  X—;  d  M  [        SU	-  5      e   [        US	 5      nU Vs0 s H&  oÝ[        [        UR                  5       5      U5      _M(     nnU Vs/ s H  ofS
   PM	     nn[        US 5      nU V	s0 s H  o™UR                  [        Xé5      5      _M     nn	X[XÏ4$ s  snf s  snf s  sn
n	f s  snf s  snf s  sn	f )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                ó   • US:H  $ rì   rú   )rc   ru   s     rC   r  Ú(JITFunction._pack_args.<locals>.<lambda>®  s	   € ¸3À+Ò;MrE   r   c                ó"   • [        U[        5      $ rG   )r.   r`   )rc   r#  s     rC   r  rÃ  ²  s   € ´ZÀÄ3Ô5GrE   )Úparse_optionsr«  r3   r8  r  ÚKeyErrorr   r   r)   rç   Ú
parse_attr)r@   r5  r	  Ú
bound_argsr>  r  r#  ÚsigkeysÚsigvalsra   Úvr  Ú
constexprsÚpathÚattrvalsr"  s                   rC   Ú
_pack_argsÚJITFunction._pack_argsŸ  sx  € à×'Ñ'¨Ó/ˆà#'§;¢;Ó/¢;˜a—6”6¡;ˆÐ/Ù!/Ó0¢˜A�Q”4¡ˆÐ0Ü(+¨GÔ(=Ô>Ò(=™f˜q�Q’TÑ(=ˆ	Ñ>à FÓ*ÐkÐ,kÓkÐ*Ø˜vÓ%ÐaÐ'aÓaÐ%Ø˜vÓ%ÐaÐ'aÓaÐ%ÛˆAØ×(Ñ(Õ(¨QÕ-=ÜÐSÐVWÑWÓXÐXñ ô # 7Ñ,MÓNˆ
Ù[eÓfÒ[eÐSWÔ-¬d°:×3DÑ3DÓ3FÓ.GÈÓNÒNÑ[eˆ
Ðfá"0Ó1¢.˜Q�a”D¡.ˆÐ1Ü˜hÑ(GÓHˆÙPUÓVÒPUÈ1�G×&Ñ&Ô'8¸Ó'EÓFÒFÑPUˆÐVà :Ð4Ð4ùò% 0ùÚ0ùÛ>ùò gùâ1ùâVs"    E»EÁEÃ-EÄE"Ä#$E'c               ó  • UR                  SU R                  5      =(       d    [        R                  R                  US'   [        R                  R
                  US'   [        R                  R                  5       n[        R                  R                  U5      nU R                   H  nU" U0 UD6  M     U R                  U   u  p‰p«nU" U0 UD6u  pÞn[        XžU5      nUR                  US 5      nUc3  U R                  X´XÞU5      u  nnnnU R                  UUUUUUU5      nUc  g [        5       nU R                   R#                  5        H8  u  u  nnu  nnUR                  UU5      =nU:w  d  M&  [%        SU SU SU 35      e   U(       dö  Uc   e['        U5      (       a  U" U5      n[)        U5      nUS   nUS:”  a  US   OSnUS:”  a  US   OSn[+        US	5      (       a  UR-                  5       nUR.                  " X/UR1                  5       Q76 nUR2                  " UUUUUR4                  UR6                  U[        R                  R8                  [        R                  R:                  /	UR1                  5       Q76   U$ )
NÚdebugÚinstrumentation_moderX   z1 has changed since we compiled this kernel, from z to r   r   r   Úresult)r|   rÒ  r   ÚruntimeÚcompilationrÓ  r   rº  Úget_current_deviceÚget_current_streamr²  Údevice_cachesr�  rÏ  Ú_do_compileÚobjectr>   r  r\   rr   r&   ÚhasattrrÔ  Úlaunch_metadatarç   r  ÚfunctionÚpacked_metadataÚlaunch_enter_hookÚlaunch_exit_hook) r@   r  r  r“   r	  r�  rÁ  r¬  Úkernel_cacher�  rª   r5  r¼  rÈ  r>  r  r  Úkernelr  rÌ  r"  Únot_presentr3   rc   ru   Úglobals_dictÚnewValÚ	grid_sizeÚgrid_0Úgrid_1Úgrid_2rÝ  s                                    rC   r  Ú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Ñ+ˆ
 GäÐ 0À'ÓJˆØ×!Ñ! # tÓ,ˆð ‰>Ø48·O±OÀGÐU_ØDKó5MÑ1ˆG�Y 
¨Eð ×%Ñ% c¨9°f¸jÈ'ÐSXÐZ`ÓaˆFØ‰~Øô “hˆØ.2×.CÑ.C×.IÑ.IÖ.KÑ*‰IˆT�1Ñ*˜˜\Ø&×*Ñ*¨4°Ó=Ð=�À#ÕEÜ"Ø& t fÐ,]Ð^aÐ]bÐbfÐgmÐfnÐoóqð qñ /Lö
 àÑ#Ð#Ð#Ü˜�~‰~Ù˜JÓ'�Ü˜D›	ˆIØ˜!‘WˆFØ )¨A£�T˜!’W°1ˆFØ )¨A£�T˜!’W°1ˆFÜ�v˜x×(Ñ(ØŸ™›�à$×4Ò4°TÐXÀJ×DUÑDUÓDWÒXˆOØ�JŠJ�v˜v v¨v°v·±È×H^ÑH^Ð`oÜ—}‘}×6Ñ6¼¿¹×8VÑ8VðnØYc×YjÑYjÓYlónàˆrE   c                óV   • U R                   c  U R                  $ U R                  U5      $ rG   )Ú_reprrO  )r@   rc   s     rC   r§  ÚJITFunction.reprì  s"   € Ø $§
¡
Ñ 2ˆt�}‰}ÐE¸¿
¹
À1»ÐErE   c	           	     ó(  >• U(       a  UO/ nU(       a  UO/ n[         TU ]  U5        UR                  U l        X l        X0l        X@l        Xpl        X€l        / U l	        [        U R                  R                  R                  5       5       H^  u  pšX“;   =(       d    U
R                  U;   nX”;   =(       d    U
R                  U;   nU R                  R                  [!        XšX¼5      5        M`     [#        U R$                  5      U l        S U l        XPl        X`l        U R                   Vs/ s H  oÝR                  PM     snU l        U R                   Vs/ s H!  oÝR0                  (       d  M  UR2                  PM#     snU l        / U l        g s  snf s  snf rG   )r1   r2   rM   rT   ÚversionrÛ   rÜ   rí  rÝ  r«  r’  r  r7  rç   r3   r9  r×   r   r½  rÙ  rã  rÒ  rY   Ú	arg_namesrí   rÙ   rÌ  r²  )r@   rF  rð  rÛ   rÜ   rÒ  rY   r§  rÝ  r•  rÝ   ÚdnsÚdns_oaÚprB   s                 €rC   r2   ÚJITFunction.__init__ï  s=  ø€ æ1BÑ-ÈÐÞKiÑ)GÐoqÐ&ä‰Ñ˜ÔØ—m‘mˆŒØŒØ!2ÔØ.LÔ+ØŒ
Ø.ÔàˆŒÜ! $§.¡.×";Ñ";×"BÑ"BÓ"DÖE‰HˆAØÑ(×K¨E¯J©JÐ:KÑ,KˆCØÑ8×h¸E¿J¹JÐJhÑ<hˆFØ�K‰K×Ñœ{¨1°SÓAÖBñ Fô )¨×);Ñ);Ó<ˆÔð ˆŒØŒ
Ø Œð +/¯+ª+Ó6ª+ QŸ&œ&©+Ñ6ˆŒØ*.¯+ª+ÓHª+ Q¿½›5˜1Ÿ5œ5©+ÑHˆŒð  ˆÕùò	 7ùÚHs   Ä'F
ÅFÅ*Fc           
     óô  • SS K nSS KJn  [        R                  R                  5       nUR                  U5      nUS   U R                  :w  a  [        SUS    SU R                   35      e[        [        US   5      nUS   n[        Xg5       VV	s0 s Hh  u  p‰UUR                  R                  U	5      (       a  UR                  U	5      O0[        U	[        5      (       a  SU	;   a  UR!                  U	S   5      OU	_Mj     n
nn	[        [        US   5      nUS	   n[        [        X¼5      5      nUS
   R#                  5        VV	s0 s H  u  p‰U[%        U	5      _M     nnn	US   R#                  5        VV	s0 s H(  u  p‰U[        U	[&        5      (       a  [        U	5      OU	_M*     nnn	US   nU R(                  U   u      nnnUR+                  U5      nU R-                  UUUU
UUSS9$ s  sn	nf s  sn	nf s  sn	nf )Nr   r3   zSpecialization data is for z but trying to preload for r  r  rk   r  r  r  r  r  T)r  )r  Útriton.languagerË   r   rº  r×  ÚloadsrO  r\   r  rŒ  r8  rÓ   Úis_dtyper.   rd  rk   r  r“  r)   rÙ  rÅ  rÚ  )r@   r¥  r  Útlr�  Údeserialized_objr  r  r  rŒ   rÌ  r  r  r"  r  r  rc   r5  s                     rC   ÚpreloadÚ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ˆô
 " -Ô?ô	
ò @‘
�ð Ø!Ÿx™x×0Ñ0°×7Ñ7ˆB�H‰H�UŒOÜ0:¸5Ä$×0GÑ0GÈKÐ[`ÓL`ˆB�L‰L˜˜{Ñ+Ô,Ðfkòlñ @ð	 	ñ 
ô œÐ 0°Ñ >Ó?ˆ
Ø% lÑ3ˆ
Ü”S˜Ó0Ó1ˆð M]Ð]hÑLi×LoÑLoÔLqÔrÒLq¹j¸c�SÔ2°5Ó9Ò9ÑLqˆ	Ñrð /¨yÑ9×?Ñ?ÔAô
âA‘
�ð ¤¨E´4×!8Ñ!8”�u”¸eÒCÙAð 	ñ 
ð ˜uÑ%ˆØ"×0Ñ0°Ñ8Ñˆˆ1ˆa�˜!Ø×'Ñ'¨Ó0ˆØ×ÑØØØØØØØð  ð 
ð 	
ùó'
ùó sùó
s   ÂA/G(ÅG.Å6/G4c           
     óH  ^ ^^^^^^^^^^^• T R                   T   u  mnmp˜T R                  [        R                  R                  TTTTTT/T5      (       a  g T R                  T TTT5      m[        R                  R                  5       n
U
bD  [        5       m[        TU	TT5      nUUU UU4S jnUUUUUUU UU4	S jnU
R                  X¼U5      nU$ T R                  TTTR                  S9nUTT'   T R                  [        R                  R                  TTTTTT/T5        U$ )Nc                 ó<   >• TR                  TTTR                  T S9$ )N)rª   r  Ú	_env_vars)r¨  r  )Úenv_varsr  r@   rA   rª   s   €€€€€rC   Úasync_compileÚ.JITFunction._do_compile.<locals>.async_compileG  s!   ø€ Ø—|‘| C°À×@PÑ@PÐ\d�|ÐeÐerE   c           
     ór   >	• U TT'   TR                  [        R                  R                  TTTTTT/T	5        g rG   )r¯  r   rÕ  Újit_post_compile_hook)
rã  r"  rÌ  r�  râ  r  r  r@   r  r  s
    €€€€€€€€€rC   Úfinalize_compileÚ1JITFunction._do_compile.<locals>.finalize_compileJ  s9   ø€ Ø$*�˜SÑ!Ø—‘¤§¡× CÑ CÀSÈ)ÐU[Ð]gÐipØ!& ¨õ1rE   )rª   r  )rÙ  r¯  r   rÕ  Újit_cache_hookr·  r   Úactive_moder|   r   r   Úsubmitr¨  r  r  )r@   r  r  r�  rÌ  r  r"  r  rc   r5  Ú
async_moder_   r  r  rã  r  râ  rA   rª   s   ````````       @@@@rC   rÚ  ÚJITFunction._do_compile:  s  ÿû€ Ø.2×.@Ñ.@ÀÑ.HÑ+ˆ�a˜ à�?‰?œ5Ÿ=™=×7Ñ7¸¸iÈÐQ[Ð]dÐglÐfmÐou×vÑvØØ�n‰n˜T 9¨j¸%Ó@ˆä#×/Ñ/×3Ñ3Ó5ˆ
ØÑ!ä6Ó8ˆHÜ% c¨7°G¸XÓFˆI÷fñ f÷1õ 1ð
  ×&Ñ& yÐAQÓRˆFð ˆð	 —\‘\ #¨f¸g×>NÑ>N�\ÐOˆFØ &ˆL˜ÑØ�O‰OœEŸM™M×?Ñ?ÀÀiÐQWÐYcÐelÐotÐnuØ"ô$àˆrE   c                ó   • [        S5      e)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)r\   ©r@   r“   r	  s      rC   Ú__call__ÚJITFunction.__call__W  s   € ÜÐWÓXÐXrE   c                óP   • SU R                    SU R                  R                   S3$ )NzJITFunction(Ú:r(  )rT   rF  r½   rI   s    rC   Ú__repr__ÚJITFunction.__repr__Z  s&   € Ø˜dŸk™k˜]¨!¨D¯G©G×,@Ñ,@Ð+AÀÐCÐCrE   )r·  r¶  rí  rñ  r¨  rÌ  rÒ  rÙ  rÛ   rÜ   rã  rÝ  rT   rY   r«  r²  rð  )r»   zbool | None)NNNNNNN)r]   rM   r½   r¾   r™  r¯  r³  r½  rÏ  r  r§  r2   rü  rÚ  r  r  rÁ   rÂ   rÃ   s   @rC   rƒ  rƒ  W  s`   ø† òð,
ð 
ô,
ò\(ò/ò5ò03òjFð mqØ;?÷" òH%
òNò:Y÷Dð DrE   rƒ  c                ó   • g rG   rú   rE  s    rC   Újitr  c  s   € àrE   ©rð  r§  rÝ  rÛ   rÜ   rÒ  rY   c                ó   • g rG   rú   r  s          rC   r  r  h  s   € ð rE   c               óF   ^^^^^^^• SUUUUUUU4S jjnU b  U" U 5      $ U$ )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                ó¬   >• [        U 5      (       d   e[        R                  R                  (       a  SSKJn  U" U TTTTTTTS9$ [        U TTTTTTTS9$ )Nr   )ÚInterpretedFunction)rð  rÛ   rÜ   rÒ  rY   r§  rÝ  )rr   r   rÕ  Ú	interpretÚinterpreterr  rƒ  )	rF  r  rÒ  rÛ   rÜ   rÝ  rY   r§  rð  s	     €€€€€€€rC   Ú	decoratorÚjit.<locals>.decorator“  sl   ø€ Ü˜�|‰|Ðˆ|Ü�=‰=×"×"Ý8Ù& r°7ÐN_ØFdÐlqØ08¸tÐUdñfð fô ØØØ"3Ø/MØØ!ØØ /ñ	ð 	rE   ©rF  r   r»   zJITFunction[T]rú   )	rF  rð  r§  rÝ  rÛ   rÜ   rÒ  rY   r  s	    ``````` rC   r  r  v  s&   þ€ ÷:õ ð& 
�~Ù˜‹}Ðð ÐrE   c                  óX   • \ rS rSrSr\S 5       rS
S jrS r\S 5       r	\S 5       r
S	rg)r  i²  zf
Can be used in place of real tensors when calling:
    kernel.warmup(MockTensor(torch.float32), ...)
c                óp   • U R                   R                  S:X  a  U R                  S:X  a  [        U 5      $ U $ )NrÓ   Útorch)rB   r]   rM   r  )r”   s    rC   r  ÚMockTensor.wrap_dtype¸  s.   € à�=‰=×!Ñ! WÓ,°·±À7Ó1JÜ˜c“?Ð"Øˆ
rE   Nc                ó(   • Uc  S/nXl         X l        g )Nr   ©rÓ   Úshape)r@   rÓ   r'  s      rC   r2   ÚMockTensor.__init__¾  s   € Ø‰=Ø�CˆEØŒ
Ø�
rE   c                óŠ   • S/nU R                   SS   H  nUR                  US   U-  5        M     [        [        U5      5      $ )Nr   rÈ   )r'  r9  rŒ  Úreversed)r@   ÚstridesÚsizes      rC   ÚstrideÚMockTensor.strideÄ  sB   € Ø�#ˆØ—J‘J˜q˜r“NˆDØ�N‰N˜7 2™;¨Ñ-Ö.ñ #ä”X˜gÓ&Ó'Ð'rE   c                 ó   • g©Nr   rú   rú   rE   rC   Údata_ptrÚMockTensor.data_ptrÊ  ó   € àrE   c                 ó   • gr0  rú   rú   rE   rC   Ú	ptr_rangeÚMockTensor.ptr_rangeÎ  r3  rE   r&  rG   )r]   rM   r½   r¾   r¿   Ústaticmethodr  r2   r-  r1  r5  rÁ   rú   rE   rC   r  r  ²  sM   † ñð
 ñó ðô
ò(ð ñó ðð ñó órE   r  c                  óT   • \ rS rSrS rS rS rSS jrS rS r	S r
S	 rS
 rS rSrg)ÚTensorWrapperiÓ  c                ó–   • X l         Xl        UR                  U l        UR                  U l        U R                  R                  U l        g rG   )rÓ   ÚbaseÚdatar�  r'  )r@   r;  rÓ   s      rC   r2   ÚTensorWrapper.__init__Õ  s1   € ØŒ
ØŒ	Ø—I‘IˆŒ	Ø—k‘kˆŒØ—Y‘Y—_‘_ˆ�
rE   c                ó6   • U R                   R                  5       $ rG   )r;  r1  rI   s    rC   r1  ÚTensorWrapper.data_ptrÜ  s   € Ø�y‰y×!Ñ!Ó#Ð#rE   c                ó4   • U R                   R                  " U6 $ rG   )r;  r-  )r@   r“   s     rC   r-  ÚTensorWrapper.strideß  s   € Ø�y‰y×Ò Ð&Ð&rE   c                ó<   • SU R                    SU R                   S3$ )NzTensorWrapper[rœ  r(  )rÓ   r;  rI   s    rC   Ú__str__ÚTensorWrapper.__str__â  s   € Ø §
¡
˜|¨2¨d¯i©i¨[¸Ð:Ð:rE   c                ó6   • U R                   R                  5       $ rG   )r;  Úelement_sizerI   s    rC   rF  ÚTensorWrapper.element_sizeå  s   € Ø�y‰y×%Ñ%Ó'Ð'rE   c                ó^   • [        U R                  R                  5       U R                  5      $ rG   )r9  r;  ÚcpurÓ   rI   s    rC   rI  ÚTensorWrapper.cpuè  s   € Ü˜TŸY™YŸ]™]›_¨d¯j©jÓ9Ð9rE   c                óN   • U R                   R                  UR                   5        g rG   )r;  Úcopy_)r@   Úothers     rC   rL  ÚTensorWrapper.copy_ë  s   € Ø�	‰	�‰˜Ÿ
™
Õ#rE   c                ó^   • [        U R                  R                  5       U R                  5      $ rG   )r9  r;  ÚclonerÓ   rI   s    rC   rP  ÚTensorWrapper.cloneî  s   € Ü˜TŸY™YŸ_™_Ó.°·
±
Ó;Ð;rE   c                ó`   • [        U R                  R                  U5      U R                  5      $ rG   )r9  r;  ÚtorÓ   )r@   r�  s     rC   rS  ÚTensorWrapper.toñ  s    € Ü˜TŸY™YŸ\™\¨&Ó1°4·:±:Ó>Ð>rE   c                ó`   • [        U R                  R                  U5      U R                  5      $ rG   )r9  r;  Ú	new_emptyrÓ   )r@   Úsizess     rC   rV  ÚTensorWrapper.new_emptyô  s"   € Ü˜TŸY™Y×0Ñ0°Ó7¸¿¹ÓDÐDrE   )r;  r<  r�  rÓ   r'  Nrù   )r]   rM   r½   r¾   r2   r1  r-  rC  rF  rI  rL  rP  rS  rV  rÁ   rú   rE   rC   r9  r9  Ó  s5   † ò%ò$ò'ô;ò(ò:ò$ò<ò?õErE   r9  c                ó
  • [        U [        5      (       a;  XR                  R                  :X  a  U R                  $ [        U R                  U5      $ [	        U S5      (       a  [        X5      $ [        S[        U 5       S35      e)Nr1  zCannot reinterpret a rD  )r.   r9  r;  rÓ   rÜ  r¯   ro   )ÚtensorrÓ   s     rC   Úreinterpretr[  ø  sm   € Ü�&œ-×(Ñ(Ø—K‘K×%Ñ%Ó%à—;‘;Ðô ! §¡¨eÓ4Ð4Ü	�˜×	$Ñ	$ä˜VÓ+Ð+äÐ/´°V³¨~¸QÐ?Ó@Ð@rE   c                óh  • U n[        U[        5      (       d#  UR                  n[        U[        5      (       d  M#  UR                  R                  R                  nUR
                  n[        UR                  5       H1  u  pEUR                  5       R                  S5      (       d  M+  X4-  n  X#4$    X#4$ )Nzdef )
r.   rZ   rF  Ú__code__Úco_filenamerL  r’  rK  rÍ   rR   )rF  Úbase_fnÚ	file_nameÚ
begin_lineÚidxÚlines         rC   Úget_jit_fn_file_linerd    s™   € Ø€GÜ˜¤+×.Ñ.Ø—*‘*ˆô ˜¤+×.Ó.à—
‘
×#Ñ#×/Ñ/€IØ×-Ñ-€Jô ˜wŸ™Ö/‰	ˆØ�:‰:‹<×"Ñ" 6×*Ó*ØÑˆJØØÐ Ð ñ	 0ð Ð Ð rE   c                  ó0   • \ rS rSrS r\S 5       rS rSrg)ÚBoundConstexprFunctioni  c                ó   • Xl         X l        g rG   )Ú__self__Ú__func__)r@   ÚinstancerF  s      rC   r2   ÚBoundConstexprFunction.__init__  s   € Ø ŒØ�rE   c                ó.   • U R                   R                  $ rG   )ri  r_   rI   s    rC   r_   Ú BoundConstexprFunction.cache_key  s   € à�}‰}×&Ñ&Ð&rE   c                óB   • U R                   " U R                  /UQ70 UD6$ rG   ©ri  rh  r  s      rC   r  ÚBoundConstexprFunction.__call__#  s   € Ø�}Š}˜TŸ]™]Ð<¨TÒ<°VÑ<Ð<rE   ro  N)	r]   rM   r½   r¾   r2   rÀ   r_   r  rÁ   rú   rE   rC   rf  rf    s    † òð ñ'ó ð'õ=rE   rf  c                  ó<   ^ • \ rS rSrU 4S jrS rSS.S jrSrU =r$ )ÚConstexprFunctioni'  c                ó$   >• [         TU ]  U5        g rG   )r1   r2   )r@   rF  rB   s     €rC   r2   ÚConstexprFunction.__init__)  s   ø€ Ü‰Ñ˜ÕrE   c                ó"   • Ub  [        X5      $ U $ rG   )rf  )r@   r$  Úobjclasss      rC   Ú__get__ÚConstexprFunction.__get__,  s   € à‰?Ü)¨#Ó4Ð4ØˆrE   N)Ú	_semanticc               ó*  • SSK JnJn  U Vs/ s H
  od" U5      PM     nnUR                  5        VVs0 s H  u  pxXt" U5      _M     nnnU R                  " U0 UD6n	Uc  U	$ [
        R                  R                  (       a  U	$ U" U	5      $ s  snf s  snnf )Nr   )Ú_unwrap_if_constexprrk   )rÊ   r{  rk   r  rF  r   rÕ  r  )
r@   ry  r“   r	  r{  rk   r#  ra   rË  Úress
             rC   r  ÚConstexprFunction.__call__2  sŒ   € ßHá15Ó6²¨AÐ$ QÖ'±ˆÐ6Ø;A¿<¹<¼>ÔJº>±°!�!Ð)¨!Ó,Ò,¹>ˆÑJð �gŠg�tÐ&˜vÑ&ˆàÑàˆJô �=‰=×"×"ØˆJÙ˜‹~Ðùò 7ùÛJs
   �B
³Brú   )	r]   rM   r½   r¾   r2   rw  r  rÁ   rÂ   rÃ   s   @rC   rr  rr  '  s   ø† õòð )-÷ ó rE   rr  c                ó   • [        U 5      $ )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.
)rr  rE  s    rC   Úconstexpr_functionr  E  s   € ô ˜RÓ Ð rE   rù   )Fr   )r§  úOptional[Callable]rÝ  r€  rÛ   úOptional[Iterable[int | str]]rÜ   r�  rÒ  úOptional[bool]rY   r‚  r»   zCallable[[T], JITFunction[T]]rG   )rF  zOptional[T]r§  r€  rÝ  r€  rÛ   r�  rÜ   r�  rÒ  r‚  rY   r‚  r»   zKernelInterface[T])FÚ
__future__r   r   r€   r/   r4   rO   rž   rP  rU  rS  Úcollectionsr   Údataclassesr   Ú	functoolsr   Útypingr   r	   r
   r   r   r   r   r   r   Útriton.backendsr   Útypesr   rN   r   r   r   Ú_utilsr   r   r   r   Úcacher   Útriton._C.libtritonr   r   r<   r;   r   ÚNodeVisitorr!   rÏ   r×   rþ   r   r&  rB  rG  rZ   r‚  r�  r“  rƒ  r  r  r9  r[  rd  rf  rr  r  rú   rE   rC   Ú<module>rŽ     s0  ðß ,Û 
Û Û Û Û Û Û 	Û Ý #Ý !Ý %ß ]× ]Õ ]å 'Ý Ý Ý Ý ß `Ó `Ý  ß Wà!€Ø3€áˆCƒL€ôg!˜Ÿ™ô g!ô^D÷4/>ñ />ôdTôZ�g˜a‘jô Zò&ò"9*òx0÷b1ñ b1ðJ ÷ð ó ðòò0	ôDD�+˜¨qÑ1ô DDðX 
óó 
ðð 
ð Ø#Ø*.Ø7;ØDHØ Ø#ñ
ð ð
ð (ð	
ð
 5ð
ð %Bð
ð ð
ð ð
ð #ô
ó 
ð
ð ð4ð Ø#Ø*.Ø7;ØDHØ Ø#ñ4Øð4ð ð	4ð
 (ð4ð 5ð4ð %Bð4ð ð4ð ð4ð ö4÷xñ ÷B"Eñ "EòJAò!ô$=˜[ô =ô˜ô ó<!rE   