
    h              	         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
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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5      *  r*\" S5      r+ " S S\RX                  5      r-S)S jr. " S S5      r/0 r0/ r1S r2S*S jr3 " S S\\+   5      r4S r5S r6S r7\ " S S5      5       r8 " S S\4\+   5      r9\S+S j5       r:\SSSSSSSS .             S,S! 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' r=S( r>g)/    )annotationsdivisionN)defaultdict)	dataclass)cached_property)
CallableGenericIterableOptionalTypeVarUnionoverloadDictAnyTupleTensorDescriptor)
ModuleType   )knobs)driver)find_paths_ifget_iterable_pathtype_canonicalisation_dictcanonicalize_dtypez.runtime.jitTc                     ^  \ rS rSrSrSU 4S jjr\S 5       rS rS 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	        0 U l
        SU l        g )Nutf-8>
   intlenmaxminlistfloatprintrangegetattr
isinstanceF)super__init__namehashlibsha256encodehasherglobals	nonlocalssupported_python_builtinsused_global_valsvisiting_arg_default_value)selfr.   r3   r4   src	__class__s        L/var/www/fran/franai/venv/lib/python3.13/site-packages/triton/runtime/jit.pyr-   DependenciesFinder.__init__)   sV    	nnSZZ%89 "*
&. TV*/'    c                6    U R                   R                  5       $ N)r2   	hexdigestr8   s    r;   retDependenciesFinder.retN   s    {{$$&&r=   c                    [         R                  " UR                  5      (       a  g[        USS5      nUR	                  [
        5      $ )NT
__module__ )inspect	isbuiltinfuncr*   
startswithTRITON_MODULE)r8   noderI   modules       r;   _is_triton_builtin%DependenciesFinder._is_triton_builtinR   s9    TYY''|R0  //r=   c                F   [        U[        5      (       Ga  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 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+   JITFunctionr6   keysRuntimeErrorr.   __name__update	cache_keystrr*   r2   r1   )r8   rI   kvar_name_v1v2func_keys           r;   _update_hashDependenciesFinder._update_hashX   s*   dK(( **//1D4I4I4N4N4PP--a0--a08&*8*KtCSTXT]T]S^^qrvrr  rA  AX  Y[  X\  \S  T  Q !!(()>)>?~~HGD*e<==HKKxw78 )r=   c                :  ^  [        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b  T R                  (       d  [        U5      [        La{  [        U[        5      (       df  [        USS5      (       dT  UR                  T R                  ;  a:  [        R                  " U5      U4T R                  UR                  [	        U5      4'   T R                  U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)r3   getr4   )r.   valr8   s     r;   name_lookup2DependenciesFinder.visit_Name.<locals>.name_lookupr   sZ    ,,""4.CDLL((..$$T40CDNN**r=   __triton_builtin__F)typectxastStoreidlocal_namesr7   r   r+   rS   r*   r5   copyr6   r`   )r8   rL   rf   re   var_dicts   `    r;   
visit_NameDependenciesFinder.visit_Namej   s    >SYY&77N77d&&&	 $DGG,
 O 77 IZ/ #344WSJ^`e=f=fGG4#A#AA>Biinh=WD!!477BxL"9:#
r=   c                b    UR                    Vs/ s H  o R                  U5      PM     sn$ s  snf r?   )eltsvisit)r8   rL   elts      r;   visit_TupleDependenciesFinder.visit_Tuple   s&     ,09959C

39555s   ,c                p   U R                  UR                  5      n[        U[        R                  5      (       a<  U R                  UR                  5      n[        U[        R                  5      (       a  M<  Ub  [        USS5      [        :X  a  g [        X!R                  5      nU R                  U5        U$ )NrV   rF   )	ru   valuer+   rk   	Attributer*   rK   attrr`   )r8   rL   lhsrB   s       r;   visit_Attribute"DependenciesFinder.visit_Attribute   s    jj$cmm,,**SYY'C cmm,,;73
B7=Hc99%#
r=   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 r?   )argsargrn   generic_visit)r8   rL   r   s      r;   visit_FunctionDef$DependenciesFinder.visit_FunctionDef   s6    /3yy~~>~GG~>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)r7   ru   )defaultsexprr8   s     r;   visit_defaults:DependenciesFinder.visit_arguments.<locals>.visit_defaults   sN    8::::26/$D'

4( % 38/%/s    A A 	A)
	itertoolschainposonlyargsr   vararg
kwonlyargsru   kw_defaultskwargr   )r8   rL   r   r   s   `   r;   visit_arguments"DependenciesFinder.visit_arguments   s    	8 ??4#3#3TYYQUQ\Q\bdfjfufuvCJJsO w 	t''(::!JJtzz"t}}%r=   c                    U R                  U5      n[        U[        5      (       a  U =R                  [	        U5      -  sl        g U R                  R                  U5        g r?   )ru   r+   r&   rn   setadd)r8   rL   targets      r;   visitAssnTarget"DependenciesFinder.visitAssnTarget   sH     D!fd##F+  (r=   c                    [        UR                  5      S:w  a  [        S5      eU R                  UR                  S   5        U R	                  U5        g )N   z2Simultaneous multiple assignment is not supported.r   )r#   targets	TypeErrorr   r   r8   rL   s     r;   visit_AssignDependenciesFinder.visit_Assign   sG    t||!
 PQQT\\!_- 	4 r=   c                \    U R                  UR                  5        U R                  U5        g r?   r   r   r   r   s     r;   visit_AnnAssign"DependenciesFinder.visit_AnnAssign   $    T[[) 	4 r=   c                \    U R                  UR                  5        U R                  U5        g r?   r   r   s     r;   	visit_ForDependenciesFinder.visit_For   r   r=   )r3   r2   rn   r.   r4   r5   r6   r7   )returnNone)rV   rE   __qualname____firstlineno____doc__r-   propertyrB   rN   r`   rq   rw   r~   r   r   r   r   r   r   __static_attributes____classcell__r:   s   @r;   r   r      s`    	#0J ' '09$%N6
!
&@)!!! !r=   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._trF   )triton.language.corelanguagecorer+   rY   striprJ   removeprefix_normalize_tyendswithpointer_type
element_tydtyper.   ri   rV   r   rd   replace)tyr   s     r;   r   r      s[   ''"cXXZ=="")Br"B==%%%%"QR&= ;;sr#2w///==r!"v...== !788  	B))	*	*=/011	B

	#	#WW	B		[[W%))"**T2*>CCr=   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 r?   )num_paramdo_not_specializedo_not_specialize_on_alignment)r8   r   paramr   r   s        r;   r-   KernelParam.__init__
  s    !2.L+r=   c                .    U R                   R                  $ r?   )r   r.   rA   s    r;   r.   KernelParam.name  s    {{r=   c                    U R                   R                  (       a2  U R                   R                  [        R                  R                  :X  a  g[        U R                   R                  5      $ )NrF   )r   
annotationrG   	Parameteremptyr   rA   s    r;   r   KernelParam.annotation  sD    {{%%)?)?7CTCTCZCZ)ZT[[3344r=   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   rF   )r   rJ   r   r   values)r8   as     r;   annotation_typeKernelParam.annotation_type  sc    OO<<!"A\\#!"A.55788??"r=   c                     SU R                   ;   $ N	constexpr)r   rA   s    r;   is_constexprKernelParam.is_constexpr&  s    doo--r=   c                    U R                   (       a  gSU R                  ;   =(       d    U R                  R                  S5      $ )NFr   r   )r   r   rJ   rA   s    r;   is_constKernelParam.is_const*  s1    $//)MT__-G-G-MMr=   c                .    U R                   R                  $ r?   )r   defaultrA   s    r;   r   KernelParam.default0  s    {{"""r=   c                d    U R                   R                  [        R                  R                  :g  $ r?   )r   r   rG   r   r   rA   s    r;   has_defaultKernelParam.has_default4  s#    {{""g&7&7&=&===r=   )r   r   r   r   N)r   r"   r   zinspect.Parameterr   boolr   r   r   rY   )rV   rE   r   r   r   r-   r   r.   r   r   r   r   r   r   r   r    r=   r;   r   r     s    LM15M     5 5
   . . N N
 # # > >r=   r   c                <   ^ ^^^ SSK Jm  SSKJm  SUUU U4S jjmT$ )Nr   r   r   r   c                ~  >^  T c  g[        T [        5      (       a  g[        T [        5      (       aC  U(       a  T" T SUS9OS nT S:X  a  U(       a  gST ::  a
  T S::  a  S	U4$ S
T ::  a
  T S::  a  SU4$ SU4$ [        T [        5      (       a  g[	        T S5      (       aa  T R
                  U4n[        R                  US 5      nUc&  US   (       a  SOS[        US   5      -   nU[        U'   U(       a  T" T SUS9OS nXd4$ [        T [        5      (       a  ST R                  4$ [        T T5      (       a  ST 4$ [	        T S5      (       a  g[        T [        5      (       aW  T  Vs/ s H  nT" U5      PM     nnU 4S jn	U	" U Vs/ s H  owS   PM	     sn5      n
U	" U Vs/ s H  owS   PM	     sn5      nX4$ [        T [        5      (       aY  [	        T R                  S5      (       d   e[        T R                  R
                  5      nSU [        T R                  5       S3S 4$ [        T T5      (       ag  [	        T R                  S5      (       d   e[        T R                  R
                  5      nSU [        T R                  5       ST R                   < S3S 4$ [#        S[%        T 5      -  5      es  snf s  snf s  snf )N)r   N)u1Nr"   )alignr   )r   r   i   ii32l            l    u64i64)fp32Ndata_ptrr   r   r   tensorr   tma_desc_cpu_ptr)	nvTmaDescNc                X   > [        TS5      (       a  [        T5      " U 6 $ [        U 5      $ )N_fields)hasattrri   tuple)valsr   s    r;   <lambda>Acreate_specialize_impl.<locals>.specialize_impl.<locals>.<lambda>d  s)    '#y:Q:Qd3i&6&bW\]aWb&br=   ztensordesc<>,zUnsupported type: %s)r+   r   r"   r'   r   r   	dtype2strrd   r   rS   rX   r   r   baser&   block_shapelayoutr   ri   )r   r   specialize_valuer   keydskresxspec
make_tupletysrT   innerGluonTensorDescriptorr   specialize_extraspecialize_impls   `            r;   r  /create_specialize_impl.<locals>.specialize_implB  s   ;&T""S!!?O"3U;UYCax,'SSI%5s|###"2s|#s|#U##!S*%%99h'C--T*C{"1vt32DSV2LL!$	#BR"3>X\C:[))//Y''%%S,--&U##0341OA&D4bJD1DqdD12CT2TtT23D;-..388Z0000&sxx~~6E!%coo)>(?qA4HH233388Z0000&sxx~~6E!%coo)>(?qaPRVWW2T#Y>?? 512s    J0 J5J:)FTT)r   r   'triton.experimental.gluon.nvidia.hopperr   )r  r  r   r  s   `@@@r;   create_specialize_implr  =  s    $a/@ /@b r=   c                    [        [        5      S:X  a  [        R                  [        S 5      5        [        S   nU" XS9S   $ )Nr   c                    g r?   r   )r\   kwargss     r;   r   mangle_type.<locals>.<lambda>x  s    PTr=   )r  )r#   specialize_impl_cacheappendr  )r   
specializer  s      r;   mangle_typer  v  s?    
 !Q&$$%;<T%UV+A.O3<Q??r=   c                  *    \ rS rSr% S\S'   SS jrSrg)KernelInterfacei}  r   runc                   ^ ^ 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$ )NFgridwarmup)r  )r   r  r"  r8   s     r;   r   -KernelInterface.__getitem__.<locals>.<lambda>  s    txx$T%'YRX'Yr=   r   )r8   r"  s   ``r;   __getitem__KernelInterface.__getitem__  s     ZYr=   r   N)r   r   )rV   rE   r   r   __annotations__r%  r   r   r=   r;   r  r  }  s    	
FZr=   r  c           
        UR                  5        VVs0 s H,  u  pVXVR                  R                  S:X  a  [        U5      OU_M.     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   r   )r.   	signatureconstant_keysconstant_vals
attrs_keys
attrs_valsoptionsr  )
itemsr:   rV   rY   jsonrT   r&   r   __dict__dumps)r.   r)  	constantsattrsr.  r  rz   r0  r	  objserialized_objs              r;   serialize_specialization_datar7    s    enetetevwevWaWZOO$<$<$Gc%jURevIwQZQ_Q_Qa?bQaAQQa?bY %**,0O,Qa,0O_cdidpdpdr_s##CC
 ZZ_N x @c0Os   3C3!C9"C>c                t   [        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                  [        [        XR                  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[(        US'   [+        UR,                  5      US'   [/        X5        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(, r   Nr   )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   rG   r   r   )r	  s    r;   r   0create_function_from_signature.<locals>.<lambda>  sC    AaDLLG,=,=,C,CCAaDaAaD6QZ[\]^[_Z`Iaar=   z
def dynamic_func(z	**optionsz):
    params = {'z': z}
    specialization = [r   z-]
    return params, specialization, options
default_rS   r  dynamic_func)r#   
parametersziprT   r   r  r   r   r   r   r+   rY   joinr&   mapr/  r   rG   r   r   rS   r  get_arg_specializationexec)sigkparamsbackendspecializationr.   kpr   r  r   rB   r   	func_bodyr   func_namespaces                 r;   create_function_from_signaturerQ    s    s~~#g,...N++-w7??!!N4&":;!#v'H$&$8$8fJ!@@GfE$TF"XJbBugQOC!!b00#66))T1R5G5G5K|5[%*
"))Br/A/A.B&T*RS #))Br/A/A.B(*KL%%/' 8, bC))DS..*>*>*@!ABk]RST U		3>>;N;N;PQ;P4QtfCv.;PQRS Txx/0 1I >>//11KD== 1 1 7 77 	)(4&5==(1   %0N=!(>w?]?](^N$% 	# .))% R
s   +J/6-J4'J4c                8    U R                    SU R                   3$ )N.)rE   r   fns    r;   get_full_namerV    s    mm_Aboo.//r=   c                  4    \ rS rSr% S\S'   S\S'   S\S'   Srg	)
JitFunctionInfoi  r   rM   rY   r.   rS   jit_functionr   N)rV   rE   r   r   r'  r   r   r=   r;   rX  rX    s    
Ir=   rX  c                     ^  \ rS rSrS r  SS jrS rS rS rS r	  SS jr
S	 r\S
 5       r\S 5       rS rS rS rS rU 4S jrU 4S jrS rSrU =r$ )rS   i  c                    g)NFr   rA   s    r;   is_gluonJITFunction.is_gluon  s    r=   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  pU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[        XXWS   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[        XU 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=](r9  r   )r)  devicer3  	num_warpsnum_ctas
num_stagesenable_fp_fusionlaunch_cooperative_gridextern_libsconfigsspecialization_data	is_warmupr  F)r  reprrU  compileis_manual_warmupalready_compiled)rU  r   rE   rF  rE  paramsr.   ra  rb  rc  rd  re  rV  r7  rf  rX  )r8   hookr  r)  r`  r3  r.  rg  ri  r.   rM   r   r   	arg_reprsrj  	full_namerh  r  s                     r;   
_call_hookJITFunction._call_hook  s    ww####IIc$++WZ[\W]F^_F^%**Rt4F^_`	{7#4#4"5[AQAQ@RR_`g`r`r_s  tG  HO  H`  H`  Ga  a{  |C  |[  |[  {\  \^  _h  ^i  ij  k!$''*	;IR[ef]gipv #" **((!,, ' 8 8'.'F'F"..#6"
 vT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)callablepre_run_hooksr  )r8   ro  s     r;   add_pre_run_hookJITFunction.add_pre_run_hook  s&    
 ~~~!!$'r=   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 XVU4$ )z!
Precompute as much as possible.
r   )CompiledKernelrk  	ASTSourcemake_backend)compilerrz  rk  r{  r|  r   activeget_current_targetrQ  r)  rn  )r8   rz  rk  r{  r|  r   rL  binders           r;   create_binderJITFunction.create_binder  sZ     	PO113v&,"/WU6F**r=   c               Z   UR                  SU R                  5      =(       d    [        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  ppU" U0 UD6u  pn[        U5      [        U5      -   nUR                  US 5      nUGc  U
R                  U5      nU R                   Vs/ s H  nUR                  PM     nnU Vs/ s H  nUS   PM
     nn[        UU5       VVs0 s H	  u  nnUU_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UUR                  ;  d  M  UU;  d  M  [!        S	U-  5      e   [#        US
 5      nU Vs0 s H'  nU[%        ['        UR)                  5       5      U5      _M)     nnU Vs/ s H  nUS   PM
     nn[#        US 5      nU Vs0 s H  nUU
R+                  [%        UU5      5      _M!     nnU R-                  [        R                  R.                  UUUUUU/U5      (       a  g U R1                  U UUU5      nU R3                  UXR                  S9nUX'   U R-                  [        R                  R4                  UUUUUU/U5        [7        5       nU R8                  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[A        U5      n"US   n#U"S:  a  US   OSn$U"S:  a  US   OSn%URB                  " X/UR)                  5       Q76 n&URD                  " U#U$U%UURF                  URH                  U&[        R                  RJ                  [        R                  RL                  /	UR)                  5       Q76   U$ 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 )Ndebugr   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   )r\   re   s     r;   r   !JITFunction.run.<locals>.<lambda>G  s	    sk?Qr=   r   c                "    [        U[        5      $ r?   )r+   rY   )r\   r	  s     r;   r   r  K  s    As9Kr=   )r   r.  rQ   z1 has changed since we compiled this kernel, from z to r   )'rd   r  r   runtimer   r~  get_current_deviceget_current_streamrv  device_cachesrY   parse_optionsrn  r.   rE  r1  KeyErrorr   r   r&   r   
parse_attrrr  jit_cache_hookr{  rk  jit_post_compile_hookobjectr6   r/  rU   ru  r#   launch_metadatar  functionpacked_metadatalaunch_enter_hooklaunch_exit_hook)'r8   r"  r#  r   r  r`  r  ro  kernel_cacher   rL  r  
bound_argsrM  r.  r  kernelr	  sigkeyssigvalsrZ   vr)  
constexprspathattrvalsr4  r9   not_presentr.   r\   re   globals_dictnewVal	grid_sizegrid_0grid_1grid_2r  s'                                          r;   r  JITFunction.run#  s&    **Wdjj9PU]]=P=Pw 11311&9 &&D$!&! ' 150B0B60J-g /5d.Ef.E+
G .!CL0!!#t, >++F3G'+{{3{!qvv{G3%34^qt^G4,/,AB,A&1aA,AIB .o0oo.6)e+ee)6)e+ee)G,,,'1A"#WZ[#[\\  'w0QRJ_ij_iW[$ 1$z7H7H7J2KT RR_iJj&45n!nH5!(,KLETYZTYqQ**+<Xq+IJJTYEZu}};;S)VU_ahkpjq%' '..y*eDC\\#f>N>N\OF &LOOEMM??iQWYcelotnu"$ h.2.C.C.I.I.K*IT1*\&**4==#E"&tf,]^a]bbfgmfnoq q /L
 ###~~J'D	I!WF )AT!W1F )AT!W1F$44TXJDUDUDWXOJJvvvvvH^H^`o}}668V8VnYcYjYjYln_ 44B k5Zs$   PP>P8.P,P#&P(c                V    U R                   c  U R                  $ U R                  U5      $ r?   )_repr_fn_name)r8   r\   s     r;   rj  JITFunction.reprm  s"     $

 2t}}E

1Er=   c	           	     P   U(       a  UO/ nU(       a  UO/ nXl         UR                  U l        X l        [        R
                  " U5      U l        X0l        X@l        [        R                  " U5      S   U l	        Xpl
        [        U5      U l        Xl        / U l        [        U R
                  R                   R#                  5       5       H^  u  pX;   =(       d    U
R$                  U;   nX;   =(       d    U
R$                  U;   nU R                  R'                  [)        XX5      5        M`     [*        R,                  " [        R.                  " U5      5      nU[0        R2                  " SU[0        R4                  5      R7                  5       S  nU R9                  U5        [;        U R<                  5      U l        S U l         0 U l!        S U l"        XPl#        X`l$        U R                   Vs/ s H  oR$                  PM     snU l%        U R                   Vs/ s H!  oRL                  (       d  M  URN                  PM#     snU l(        / U l)        URT                  U l*        URV                  U l+        URX                  U l,        URZ                  U l-        UR                  U l        g s  snf s  snf )Nr   z^def\s+\w+\s*\().rU  rE   rM   versionrG   r)  r   r   getsourcelinesstarting_line_numberr  rV  r  r  rn  	enumeraterD  r   r.   r  r   textwrapdedent	getsourceresearch	MULTILINEstart_unsafe_update_srcr   r  r  hashr6   r  r  rR   	arg_namesr   r   r  rv  r   rV   r   __globals__)r8   rU  r  r   r   r  rR   rj  r  ir   dnsdns_oar9   ps                  r;   r-   JITFunction.__init__p  s   1B-Ki)Goq&mm **2.!2.L+$+$:$:2$>q$A!
%b).!$..";";"B"B"DEHA(KEJJ:K,KC8hEJJJh<hFKK{1SAB F oog//34")).R\\BHHJKL$(););<	 TV 
  +/++6+Q&&+6*.++H+Q5155+H   zzOO>>-- 7Hs   &JJ#)J#c                p    U R                   [        R                  " U R                  5      R                  -  $ r?   )r  rG   getclosurevarsrU  r4   rA   s    r;   get_capture_scopeJITFunction.get_capture_scope  s(    '"8"8"A"K"KKKr=   c                   U R                   c  [        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        U R                   $ )N)r.   r3   r4   r9   )r  rG   r  rU  r4   r   r  r  r9   ru   parserB   rY   r  dictsortedr6   r/  )r8   r4   dependencies_finders      r;   rX   JITFunction.cache_key  s     99..tww7AAI"4$--QUQaQamv9=#C%%djjl3+//#d6O6O2PPDI$(0C0T0T0Z0Z0\)]$^D!yyr=   c                    SSK Jn  U$ )Nr   r   )r   r   )r8   r   s     r;   ri   JITFunction.type  s    2r=   c               \    U R                   " [        [        R                  U5      USS.UD6$ )NTr!  )r  rG  
MockTensor
wrap_dtype)r8   r"  r   r  s       r;   r#  JITFunction.warmup  s(    xxZ5J5JD1QT$\U[\\r=   c           	     N   SSK JnJn  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	[        X5       V
Vs0 s H8  u  pXR                  R!                  U5      (       a  UR                  U5      OU_M:     nn
n[        [        US	   5      nUS
   n[#        [        X5      5      n[#        US   R%                  5       5      nU" U UX5      nUS   R%                  5        V
Vs0 s H(  u  pU
['        U[(        5      (       a  [        U5      OU_M*     nn
nUS   n
U" US U5      nUU R*                  U   S   U
'   U$ s  snn
f s  snn
f )Nr   )rk  r{  r   r.   zSpecialization data is for z but trying to preload for r*  r+  r,  r-  r)  r.  r  )r}  rk  r{  r0  triton.languager   r   r~  r  loadsr  rU   rG  r   rE  r   is_dtyper  r/  r+   r&   r  )r8   rh  rk  r{  r0  tlr`  deserialized_objr*  r+  r  rz   r3  r,  r-  r4  r)  r9   r.  r  s                       r;   preloadJITFunction.preload  s   1$113::&9:F#t}}4-.>v.F-GGbcgcpcpbqrt tE#3O#DE(9 "-?
?
 HH$5$5e$<$<%%G? 	 
  0 >?
%l3
S01)+6<<>?	i: /y9??A
A
 E4!8!8ueCA 	 
 u%dG,-36"1%c*!

s   ?F/F!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   )rk   r  r9   r+   Moduler#   bodyFunctionDef)r8   trees     r;   r  JITFunction.parse  se    yy"$

++++499~"""$))A,8888r=   c                    [        S5      e)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rU   )r8   r   r  s      r;   __call__JITFunction.__call__  s    WXXr=   c                V   > US:X  a  [        SU S35      e[        [        U ]  X5        g )Nr9   zCannot set attribute 'zX' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)AttributeErrorr,   rS   __setattr__)r8   r.   rz   r:   s      r;   r  JITFunction.__setattr__  s:    5= #9$ @, "- . . 	k4,T9r=   c                4   > SU l         [        TU ]	  SU5        g)zv
The only method allowed to modify src.
Bypasses the __setattr__ restriction by calling super().__setattr__ directly.
Nr9   )r  r,   r  )r8   new_srcr:   s     r;   r  JITFunction._unsafe_update_src  s    
 	E7+r=   c                P    SU R                    SU R                  R                   S3$ )NzJITFunction(:r9  )rM   rU  r   rA   s    r;   __repr__JITFunction.__repr__  s&    dkk]!DGG,@,@+ACCr=   )r{  rz  r   r  rE   rV   r   r  r  r  rk  r  r  r  r   r   rU  r  r  r  rM   rR   rn  rv  r)  r  r6   r  )r   zbool | None)NNNNNNN)rV   rE   r   r   r\  rr  rw  r  r  rj  r-   r  r   rX   ri   r#  r  r  r  r  r  r  r   r   r   s   @r;   rS   rS     s    ,
 
,
\(+HTF mq;?<(|L 	 	  ]@Y:,D Dr=   rS   c                    g r?   r   rT  s    r;   jitr    s    r=   r  rj  r  r   r   r  rR   c                    g r?   r   r  s          r;   r  r    s     r=   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  rR   rj  r  )ru  r   r  	interpretinterpreterr  rS   )	rU  r  r  r   r   r  rR   rj  r  s	     r;   	decoratorjit.<locals>.decorator8  sl    |||==""8&r7N_Fdlq08tUdf f "3/M! /	 	r=   rU  r   r   zJITFunction[T]r   )	rU  r  rj  r  r   r   r  rR   r  s	    ``````` r;   r  r    s&    : & 
~} r=   c                  N    \ rS rSrSr\S 5       rS r\S 5       r\S 5       r	Sr
g)	r  iW  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)r:   rV   rE   r  )r   s    r;   r  MockTensor.wrap_dtype]  s.    ==!!W,71Jc?"
r=   c                    Xl         g r?   r   )r8   r   s     r;   r-   MockTensor.__init__c  s    
r=   c                     gNr   r   r   r=   r;   r   MockTensor.data_ptrf      r=   c                     gr  r   r   r=   r;   	ptr_rangeMockTensor.ptr_rangej  r  r=   r   N)rV   rE   r   r   r   staticmethodr  r-   r   r  r   r   r=   r;   r  r  W  sH    
  
    r=   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)TensorWrapperio  c                    X l         Xl        UR                  U l        UR                  U l        U R                  R                  U l        g r?   )r   r  datar`  shape)r8   r  r   s      r;   r-   TensorWrapper.__init__q  s1    
	II	kkYY__
r=   c                6    U R                   R                  5       $ r?   )r  r   rA   s    r;   r   TensorWrapper.data_ptrx  s    yy!!##r=   c                4    U R                   R                  " U6 $ r?   )r  stride)r8   r   s     r;   r  TensorWrapper.stride{  s    yy&&r=   c                <    SU R                    SU R                   S3$ )NzTensorWrapper[r_  r9  )r   r  rA   s    r;   __str__TensorWrapper.__str__~  s    

|2dii[::r=   c                6    U R                   R                  5       $ r?   )r  element_sizerA   s    r;   r  TensorWrapper.element_size  s    yy%%''r=   c                ^    [        U R                  R                  5       U R                  5      $ r?   )r  r  cpur   rA   s    r;   r  TensorWrapper.cpu  s    TYY]]_djj99r=   c                N    U R                   R                  UR                   5        g r?   )r  copy_)r8   others     r;   r  TensorWrapper.copy_  s    		

#r=   c                ^    [        U R                  R                  5       U R                  5      $ r?   )r  r  cloner   rA   s    r;   r#  TensorWrapper.clone  s    TYY__.

;;r=   c                `    [        U R                  R                  U5      U R                  5      $ r?   )r  r  tor   )r8   r`  s     r;   r&  TensorWrapper.to  s     TYY\\&14::>>r=   c                `    [        U R                  R                  U5      U R                  5      $ r?   )r  r  	new_emptyr   )r8   sizess     r;   r)  TensorWrapper.new_empty  s"    TYY007DDr=   )r  r  r`  r   r  Nr   )rV   rE   r   r   r-   r   r  r  r  r  r  r#  r&  r)  r   r   r=   r;   r  r  o  s5    %$';(:$<?Er=   r  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)Nr   zCannot reinterpret a rS  )r+   r  r  r   r   r   ri   )r   r   s     r;   reinterpretr-    sm    &-((KK%%%;; !e44		$	$V++/V~Q?@@r=   c                   U n[        U[        5      (       d#  UR                  n[        U[        5      (       d  M#  UR                  R                  R                  n[
        R                  " UR                  5      u  p4[        U5       H1  u  pVUR                  5       R                  S5      (       d  M+  XE-  n  X$4$    X$4$ )Nzdef )
r+   rS   rU  __code__co_filenamerG   r  r  r   rJ   )rU  base_fn	file_namelines
begin_lineidxlines          r;   get_jit_fn_file_liner7    s    G+..** +..

##//I..wzz:E u%	::<""6**J  	 &   r=   r   )Fr  )rj  Optional[Callable]r  r8  r   Optional[Iterable[int | str]]r   r9  r  Optional[bool]rR   r:  r   zCallable[[T], JITFunction[T]]r?   )rU  zOptional[T]rj  r8  r  r8  r   r9  r   r9  r  r:  rR   r:  r   z4Union[JITFunction[T], Callable[[T], JITFunction[T]]])?
__future__r   r   rk   ro   r/   rG   r   r  r  collectionsr   dataclassesr   	functoolsr   typingr   r	   r
   r   r   r   r   r   r   r   triton.tools.tensor_descriptorr   typesr   rF   r   runtime.driverr   _utilsr   r   r   r   rV   r#   rK   r   NodeVisitorr   r   r   r  r  r  r  r  r7  rQ  rV  rX  rS   r  r  r  r-  r7  r   r=   r;   <module>rE     s   , 
     	  # ! % d d d ;   # e e.3~../CLH! H!`D4/> />d 	 6r@	Zgaj 	Z	7*t0   eD/!$ eDZ	 
 
 
 #*.7;DH #
 
 (	

 5
 %B
 
 
 #
 

 4 #*.7;DH #44 	4
 (4 54 %B4 4 4 :4x 0"E "EJA!r=   