
    Ji|              	         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        T/var/www/html/dynamic-report/venv/lib/python3.13/site-packages/triton/runtime/jit.pyr2   DependenciesFinder.__init__.   sn    	nnSZZ%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    TYY''|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&   $,,,, &&++-0E0E0J0J0LLAKH))!,EB))!,EBx"&xjB4?OPTPYPY{Zmnrn{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!6z#?Y?Y!B3%HII **;?==;Mx:XD!!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_lookup2DependenciesFinder.visit_Name.<locals>.name_lookup   sZ    ,,""4.CDLL((..$$T40CDNN**rE   )ro   ctxastStorert   local_namesr:   rq   )r@   rS   r}   ru   rv   s   `    rC   
visit_NameDependenciesFinder.visit_Name   s|    >SYY&77N77d&&&	 $DGG,77d444JcTWW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&     ,09959C

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    jj$cmm,,**SYY'C cmm,,3
B/;(&<&<<c99%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    /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)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#3TYYQUQ\Q\bdfjfufuvCJJsO w 	t''(::!JJtzz"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!fd##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    	)0V ' '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XXZ=="")Br"B==%%%%"QR&= ;;sr#2w///==r!"v...== !788  	B))	*	*=/011	B

	#	#WW	B		[[W%))"**T2*>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    {{%%)?)?7CTCTCZCZ)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    OO<<!"A\\#!"A.55788??"rE   c                     SU R                   ;   $ Nrk   )r   rI   s    rC   is_constexprKernelParam.is_constexprJ  s    doo--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    $//)MT__-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!+sjPQR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(    xxZ5J5JD1QT$\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    txx$T%'YRX'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    $//+ ,JC 	??33w>SZ&+oo&>&>+&Mekk"SX	Y+   QZQ_Q_Qa?bQaAQQa?bY %**,0O,Qa,0O_cdidpdpdr_s##CC
 ZZ_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                  [        [        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[(        nXS'   X-S'   [*        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(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    AaDLLG,=,=,C,CCAaDaAaD6QZ[\]^[_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??!!N4&":;!#v'H$&$8$8fJ!@@GfE-dV2hZr*RPUwVWXC!!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   -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    mm_Aboo.//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=6L6LR6P3DL$3 &b)#//+ oobggdll34")).R\\BHHJKL		 TV zzOO>>--7  	YPQWX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    __yy$yy _
 %T]]O4DI..tww7AAI"4$--QUQaQamv9=#C%%djjl3+//#d6O6O2PPDI$(0C0T0T0Z0Z0\)]$^D!6II9=9N9N9T9T9V=9V"5)4Xc!+C!; *tk9V= > >I  tyy'7'7'@AKKMDI# $ yy	= _$ yys*   GDGG#G*AGG
G0c                ,    [        U R                  5      $ rG   )rZ  r_   rI   s    rC   __hash__JITCallable.__hash__  s    DNN##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    yy#$

++++499~"""$))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    yy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
0CrE   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     ?3C*3//3s   )r.   r)   r   rB   tuplerZ   r_   )r$  r   resultsr  s      rC   r  ,compute_cache_key.<locals>.replace_callables:  s    c4  69:cs%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EI%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>    dD!! %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  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=](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    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)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2v..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  pX_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	XR                  ;  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  oUR                  [        X5      5      _M     nn	X[X4$ 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35GrE   )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   ''/#';;/;a66;/!/0AQ40(+G(=>(=fqQT(=	>F*k,kk*v%a'aa%v%a'aa%A(((Q-=SVWWXX  #7,MN
[ef[eSW-d:3D3D3F.GNN[e
f"01.QaD.1h(GHPUVPU1G&&'8'EFFPU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  ppnU" U0 UD6u  pn[        XU5      nUR                  US 5      nUc3  U R                  XXU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    **Wdjj9PU]]=P=Pw).):):)O)O%& 11311&9 &&D$!&! ' CGBTBTU[B\? /5d.Ef.E+
G 0'J!!#t, >48OOGU_DK5M1GY
E %%c9fj'SXZ`aF~ 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vx(($44TXJDUDUDWXOJJvvvvvH^H^`o}}668V8VnYcYjYjYlnrE   c                V    U R                   c  U R                  $ U R                  U5      $ rG   )_reprrO  )r@   rc   s     rC   r  JITFunction.repr  s"     $

 2t}}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        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`     [#        U R$                  5      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!  oR0                  (       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)Goq&mm!2.L+
.!$..";";"B"B"DEHA(KEJJ:K,KC8hEJJJh<hFKK{1SAB F )););< 
  +/++6+Q&&+6*.++H+Q5155+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  pU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[        [        X5      5      nUS
   R#                  5        VV	s0 s H  u  pU[%        U	5      _M     nn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 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-GGbcgcpcpbqrt tE#3O#DE(9
 "-?	
 @
 !xx0077BHHUO0:5$0G0GK[`L`BLL{+,fkl @	 	 
  0 >?
%l3
S01 M]]hLiLoLoLqrLqjcS2599Lq	r /y9??A
A
 E4!8!8ueCA 	 
 u%"0081a!''0   
 	
'
 s
s   A/G(G.6/G4c           
     H  ^ ^^^^^^^^^^^ T R                   T   u  mnmpT 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                  XU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_compile1JITFunction._do_compile.<locals>.finalize_compileJ  s9    $*S! C CS)U[]gip!&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[]dglfmouvvnnT9j%@#//335
!68H%c7GXFIf f1 1
  &&yAQRF 	 \\#fg>N>N\OF &LOOEMM??iQWYcelotnu"$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&    dkk]!DGG,@,@+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`    ,
 
,
\(/503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&r7N_Fdlq08tUd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,71Jc?"
rE   Nc                (    Uc  S/nXl         X l        g )Nr   r   shape)r@   r   r'  s      rC   r2   MockTensor.__init__  s    =CE

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    #JJqrNDNN72;-. #Xg&''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    
	II	kkYY__
rE   c                6    U R                   R                  5       $ rG   )r;  r1  rI   s    rC   r1  TensorWrapper.data_ptr  s    yy!!##rE   c                4    U R                   R                  " U6 $ rG   )r;  r-  )r@   r   s     rC   r-  TensorWrapper.stride  s    yy&&rE   c                <    SU R                    SU R                   S3$ )NzTensorWrapper[r  r(  )r   r;  rI   s    rC   __str__TensorWrapper.__str__  s    

|2dii[::rE   c                6    U R                   R                  5       $ rG   )r;  element_sizerI   s    rC   rF  TensorWrapper.element_size  s    yy%%''rE   c                ^    [        U R                  R                  5       U R                  5      $ rG   )r9  r;  cpur   rI   s    rC   rI  TensorWrapper.cpu  s    TYY]]_djj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    TYY__.

;;rE   c                `    [        U R                  R                  U5      U R                  5      $ rG   )r9  r;  tor   )r@   r  s     rC   rS  TensorWrapper.to  s     TYY\\&14::>>rE   c                `    [        U R                  R                  U5      U R                  5      $ rG   )r9  r;  	new_emptyr   )r@   sizess     rC   rV  TensorWrapper.new_empty  s"    TYY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    &-((KK%%%;; !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    H156A$Q'6;A<<>J>!!)!,,>J ggt&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CLg! g!^D4/> />dTZgaj Z&"9*x0b1 b1J   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   