
    (ThL              	         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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Jr  \ S\!" S
5      *  r"\" S5      r# " S S\RH                  5      r%SGS jr& " S S5      r'0 r(/ r)S r*SHS jr+ " S S\\#   5      r,S r-S r.0 SS_SS_SS_SS_SS_S S!_S"S!_S#S_S$S%_S&S%_S'S(_S)S*_S+S,_S-S._S/S0_S1S2_S3S4_S5S6S7S8S9S:.Er/\0" \/Rc                  5       5       H  r2\2\/\2'   M
      " S; S<\,\#   5      r3\SIS= j5       r4\SSSSSSSS>.             SJS? jj5       r4 SKSSSSSSSS>.               SLS@ jjjr4 " SA SB5      r5 " SC SD5      r6SE r7SF r8g)M    )annotationsdivisionN)defaultdict)cached_property)
CallableGenericIterableOptionalTypeVarUnionoverloadDictAnyTuple   )driver)
ModuleType)find_paths_ifget_iterable_pathz.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        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supported_python_builtinsused_global_valsvisiting_arg_default_value)selfr(   r-   src	__class__s       J/var/www/auris/envauris/lib/python3.13/site-packages/triton/runtime/jit.pyr'   DependenciesFinder.__init__%   sQ    	nnSZZ%89 *
&. TV*/'    c                6    U R                   R                  5       $ N)r,   	hexdigestr1   s    r4   retDependenciesFinder.retI   s    {{$$&&r6   c                    [         R                  " UR                  5      (       a  g[        USS5      nUR	                  [
        5      $ )NT
__module__ )inspect	isbuiltinfuncr$   
startswithTRITON_MODULE)r1   noderB   modules       r4   _is_triton_builtin%DependenciesFinder._is_triton_builtinM   s9    TYY''|R0  //r6   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%   JITFunctionr/   keysRuntimeErrorr(   __name__update	cache_keystrr$   r,   r+   )r1   rB   kvar_name_v1v2func_keys           r4   _update_hashDependenciesFinder._update_hashS   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 )r6   c                N   [        UR                  5      [        R                  L a  UR                  $ UR                  U R
                  ;   a  g U R                  R                  UR                  S 5      nUb  U R                  (       d  [        U5      [        Laz  [        U[        5      (       de  [        USS5      (       dS  UR                  U R                  ;  a9  X R                  4U R                  UR                  [	        U R                  5      4'   U R                  U5        U$ )N__triton_builtin__F)typectxastStoreidlocal_namesr-   getr0   r   r%   rL   r$   r.   r/   rY   )r1   rE   vals      r4   
visit_NameDependenciesFinder.visit_Namee   s    >SYY&77N77d&&&lltww-
 O 77 IZ/ #344WSJ^`e=f=fGG4#A#AABE||ATD!!477Bt||,<"=>#
r6   c                b    UR                    Vs/ s H  o R                  U5      PM     sn$ s  snf r8   )eltsvisit)r1   rE   elts      r4   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$ )NrO   r?   )	ri   valuer%   r_   	Attributer$   rD   attrrY   )r1   rE   lhsr;   s       r4   visit_Attribute"DependenciesFinder.visit_Attribute   s    jj$cmm,,**SYY'C cmm,,;73
B7=Hc99%#
r6   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 r8   )argsargrb   generic_visit)r1   rE   rv   s      r4   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)r0   ri   )defaultsexprr1   s     r4   visit_defaults:DependenciesFinder.visit_arguments.<locals>.visit_defaults   sN    8::::26/$D'

4( % 38/%/s    A A 	A)
	itertoolschainposonlyargsru   vararg
kwonlyargsri   kw_defaultskwargr|   )r1   rE   r~   rv   s   `   r4   visit_arguments"DependenciesFinder.visit_arguments   s    	8 ??4#3#3TYYQUQ\Q\bdfjfufuvCJJsO w 	t''(::!JJtzz"t}}%r6   c                    U R                  U5      n[        U[        5      (       a  U =R                  [	        U5      -  sl        g U R                  R                  U5        g r8   )ri   r%   r    rb   setadd)r1   rE   targets      r4   visitAssnTarget"DependenciesFinder.visitAssnTarget   sH     D!fd##F+  (r6   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   rw   r1   rE   s     r4   visit_AssignDependenciesFinder.visit_Assign   sG    t||!
 PQQT\\!_- 	4 r6   c                \    U R                  UR                  5        U R                  U5        g r8   r   r   rw   r   s     r4   visit_AnnAssign"DependenciesFinder.visit_AnnAssign   $    T[[) 	4 r6   c                \    U R                  UR                  5        U R                  U5        g r8   r   r   s     r4   	visit_ForDependenciesFinder.visit_For   r   r6   )r-   r,   rb   r(   r.   r/   r0   )returnNone)rO   r>   __qualname____firstlineno____doc__r'   propertyr;   rG   rY   re   rk   rr   rx   r   r   r   r   r   __static_attributes____classcell__r3   s   @r4   r   r      s_    	"0H ' '09$<6
!
&@)!!! !r6   r   c                    [        U [        5      (       a  U R                  $ [        U [        5      (       a  U $ [	        U 5      $ r8   )r%   r]   rO   rR   repr)tys    r4   _normalize_tyr      s4    "d{{	B			8Or6   c                      \ rS rSrSr  SS jr\S 5       r\S 5       r\S 5       r	\S 5       r
\S 5       r\S	 5       r\S
 5       rSrg)KernelParam   zBRepresents a parameter (name plus metadata) to a @jit'ed function.c                4    Xl         X l        X0l        X@l        g r8   )num_paramdo_not_specializedo_not_specialize_on_alignment)r1   r   paramr   r   s        r4   r'   KernelParam.__init__   s    !2.L+r6   c                .    U R                   R                  $ r8   )r   r(   r:   s    r4   r(   KernelParam.name   s    {{r6   c                    U R                   R                  (       a2  U R                   R                  [        R                  R                  :X  a  g[        U R                   R                  5      $ )Nr?   )r   
annotationr@   	Parameteremptyr   r:   s    r4   r   KernelParam.annotation   sD    {{%%)?)?7CTCTCZCZ)ZT[[3344r6   c                    U R                   nS H:  u  p#XR                  U5      [        U5      -   S  nU(       d  M-  X!;   d  M4  U U 3s  $    US:X  a  gg)N))uintu)r   iboolu1r?   )r   findr   )r1   r   ty1ty2widths        r4   annotation_typeKernelParam.annotation_type   s\    __
5HCs3c#h>?@Eu*ug& 6 r6   c                     SU R                   ;   $ N	constexpr)r   r:   s    r4   is_constexprKernelParam.is_constexpr  s    doo--r6   c                P    SU R                   ;   =(       a    U R                  (       + $ )Nconst)r   r   r:   s    r4   is_constKernelParam.is_const
  s    $//)C$2C2C.CCr6   c                .    U R                   R                  $ r8   )r   defaultr:   s    r4   r   KernelParam.default  s    {{"""r6   c                d    U R                   R                  [        R                  R                  :g  $ r8   )r   r   r@   r   r   r:   s    r4   has_defaultKernelParam.has_default  s#    {{""g&7&7&=&===r6   )r   r   r   r   N)r   r   r   zinspect.Parameterr   r   r   r   )rO   r>   r   r   r   r'   r   r(   r   r   r   r   r   r   r   r    r6   r4   r   r      s    LM15M     5 5
   . . D D # # > >r6   r   c                 z   ^ ^ [         (       a	  [         S   $ SSKJm   SU U4S jjm[         R                  T5        T$ )Nr   )r   c                  >^  T c  g[        T [        5      (       a  ST R                  4$ [        T T5      (       a  ST 4$ [        T [        5      (       a  g[        T [        5      (       aC  U(       a  U" 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      (       a  g[        T [        5      (       aW  T  Vs/ s H  nT" Xa5      PM     nnU 4S jnU" U Vs/ s H  ofS   PM	     sn5      n	U" U Vs/ s H  ofS   PM	     sn5      n
X4$ T R                  U4n[        R                  US 5      nUc?  US   (       a  SOS[        [        US   5      R                  S5      S      -   nU[        U'   U(       a  U" T SUS9OS nX4$ s  snf s  snf s  snf )N)r   Nr   )i1Nr   )alignr   )r   r   i   ii32l            l    u64i64)fp32Ntma_desc_cpu_ptr)	nvTmaDescNc                X   > [        TS5      (       a  [        T5      " U 6 $ [        U 5      $ )N_fields)hasattrr]   tuple)valsrv   s    r4   <lambda>Acreate_specialize_impl.<locals>.specialize_impl.<locals>.<lambda>;  s)    '#y:Q:Qd3i&6&bW\]aWb&br6   r   z*k*.r   tensor)r%   rL   rQ   r   r   r!   r   r   dtype	dtype2strrc   type_canonicalisation_dictrR   split)rv   specialize_extrar   specialize_valuer   keyxspec
make_tupletysrM   dskresr   specialize_impls   `            r4   r   /create_specialize_impl.<locals>.specialize_impl!  s   ;&[))//Y''%%T""S!!?O"3U;UYCax,'SSI%5s|###"2s|#s|#U##!S,--&U##BEF#QOA8#DFbJD1DqdD12CT2TtT23D; 99h'C--T*C{"1vt32LSQTUVQW[M^M^_bMcdfMg2hh!$	#BR"3>X\C: G12s   .F;G )G)FTT)specialize_impl_cachelanguager   append)r   r   s   @@r4   create_specialize_implr     s8    $R(($& &P   1r6   c                .    [        5       nU" U S US9S   $ )Nc                    g r8   r   )rU   kwargss     r4   r   mangle_type.<locals>.<lambda>O  s    Dr6   )r   r   )r   )rv   
specializer   s      r4   mangle_typer  M  s     ,.O3 8:VWXYYr6   c                  *    \ rS rSr% S\S'   SS jrSrg)KernelInterfaceiR  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  )ru   r  r  r1   s     r4   r   -KernelInterface.__getitem__.<locals>.<lambda>[  s    txx$T%'YRX'Yr6   r   )r1   r  s   ``r4   __getitem__KernelInterface.__getitem__U  s     ZYr6   r   N)r   r   )rO   r>   r   r   __annotations__r  r   r   r6   r4   r  r  R  s    	
FZr6   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   )
itemsr3   rO   rR   jsonrM   r    values__dict__dumps)r(   r  	constantsattrsr  r   rn   r  r   objserialized_objs              r4   serialize_specialization_datar"  _  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                   [        U R                  5      [        U5      :X  d   e/ n[        U R                  R                  5       U5       H  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                  SUR                   S	U	 S
35        M  UR                  U	 5        M     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'   ['        5       US'   UR(                  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(z, specialize_extra, , z("z",) + z[1:]c                z    U S   R                   [        R                  R                  L a  U S   $ U S    SU S    3$ )Nr   r   z	=default_)r   r@   r   r   )r   s    r4   r   0create_function_from_signature.<locals>.<lambda>  sC    AaDLLG,=,=,C,CCAaDaAaD6QZ[\]^[_Z`Iaar6   z
def dynamic_func(z	**optionsz):
    params = {'z': z}
    specialization = [,z-]
    return params, specialization, options
default_rL   r   r   dynamic_func)r   
parametersziprM   r   r   r   r   r   r   joinr    mapr  r   r@   r   r   rL   r   get_arg_specializationexec)sigkparamsbackendspecializationr(   kpr   r  r   r;   rv   	func_bodyr   func_namespaces                 r4   create_function_from_signaturer;  k  s2    s~~#g,...N++-w7??!!N4&":;!#v'H$&$8$8fJ!@@GfE$TF*>xj:,VXY^X__`aC!!%%2+=+=*>fSE&NO%%/ 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=!(>(@N$%)0)G)GN%& 	# .))' R
s   <I-I	8I	r   r   
float8e4nvfp8e4nvfloat8e5fp8e5float8e4b15fp8e4b15float8_e4m3fn
float8e4b8fp8e4b8float8_e4m3fnuzfloat8_e5m2float8e5b16fp8e5b16float8_e5m2fnuzfloat16fp16bfloat16bf16float32r   float64fp64int8i8int16i16int32r   r   u8u16u32r   )int64uint8uint16uint32uint64c                     ^  \ rS rSrSrSrS rS rS rS r	S r
  SS j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$ )rL   i  Nc	                   U(       a  [         R                  O[         R                  n	U	c  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 " S S5      n[        XXFS   XQ5      nUUUUR                  UR                  UR                  UR                  UR                  UR                   UUUS.nU	" UUU" XU 5      SU0UEUSS9$ s  snnf )NFr'  r   z: z[num_warps=z, num_ctas=z, num_stages=z, enable_fp_fusion=z, launch_cooperative_grid=](r$  c                      \ rS rSrS rSrg)/JITFunction._call_hook.<locals>.JitFunctionInfoi  c                (    Xl         X l        X0l        g r8   )rF   r(   jit_function)r1   rF   r(   rd  s       r4   r'   8JITFunction._call_hook.<locals>.JitFunctionInfo.__init__  s    $ 	$0!r6   )rd  rF   r(   N)rO   r>   r   r   r'   r   r   r6   r4   JitFunctionInforb    s    r6   rf  r   )r  devicer  	num_warpsnum_ctas
num_stagesenable_fp_fusionlaunch_cooperative_gridextern_libsconfigsspecialization_data	is_warmupr   )r   r   fncompileis_manual_warmupalready_compiled)rL   
cache_hookcompiled_hookrq  rO   r>   r0  r/  paramsr(   rh  ri  rj  rk  rl  r"  rm  )r1   r   r  rg  r  r  rn  rp  beforehookr(   rF   r   r   	arg_reprsr   rf  ro  r  s                      r4   
_call_hookJITFunction._call_hook  s    *0{%%[5N5N<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	 	 <DY`aXbdkq #" **((!,, ' 8 8'.'F'F"..#6"
 vT2C*6*&"
 	
9 `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   )r1   ry  s     r4   add_pre_run_hookJITFunction.add_pre_run_hook  s&    
 ~~~!!$'r6   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   )CompiledKernelrr  	ASTSourcemake_backend)compilerr  rr  r  r  r   activeget_current_targetr;  r  rw  )r1   r  rr  r  r  r   r6  binders           r4   create_binderJITFunction.create_binder   sZ     	PO113v&,"/WU6F**r6   c                  UR                  SU R                  5      =(       d#    [        R                  R                  SS5      S:H  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-                  UUUUUU/USS9(       a  g U R/                  U UUU5      nU R1                  UXR                  S9nUX'   U R-                  UUUUUU/USS9  [3        5       nU R4                  R7                  5        H8  u  u  nnu  nn U R                  UU5      =n!U:w  d  M&  [9        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R>                  " X/UR)                  5       Q76 n&UR@                  " U#U$U%UURB                  URD                  U&U RF                  RH                  U RF                  RJ                  /	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debugTRITON_DEBUG01r   device_typez=device_type option is deprecated; current target will be usedrg  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   )rU   rd   s     r4   r   !JITFunction.run.<locals>.<lambda>/  s	    sk?Qr6   r   c                "    [        U[        5      $ r8   )r%   rR   )rU   r   s     r4   r   r  3  s    As9Kr6   T)rx  )r   r  FrJ   z1 has changed since we compiled this kernel, from z to r   )&rc   r  osenvironr   r  get_current_deviceget_current_streamr  device_cachesrR   parse_optionsrw  r(   r/  r  KeyErrorr   r   r    r  
parse_attrr{  r  rr  objectr/   r  rN   r~  r   launch_metadatar  functionpacked_metadatar  launch_enter_hooklaunch_exit_hook)'r1   r  r  ru   r  rg  r  ry  kernel_cacher   r6  r  
bound_argsr7  r  r   kernelr   sigkeyssigvalsrS   vr  
constexprspathattrvalsr  r2   not_presentr(   rU   rd   globals_dictnewVal	grid_sizegrid_0grid_1grid_2r  s'                                          r4   r  JITFunction.run  s    **Wdjj9gRZZ^^N\_=`dg=gw 11311&9 &&D$!&! ' 150B0B60J-g.4d.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sIvz7UGU[dhi..y*eDC\\#f>N>N\OF &LOOCFJ%RXafOg 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^&(;(;(M(MtObObOsOs-"))+- ] 44B k5Zs$   O'O O%.O+5O0&O5c                V    U R                   c  U R                  $ U R                  U5      $ r8   )_repr_fn_name)r1   rU   s     r4   r   JITFunction.reprT  s"     $

 2t}}E

1Er6   c	           	     0   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R                  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R                  U l        URV                  U l+        UR                  U l        g s  snf s  snf )Nr   z^def\s+\w+\s*\(),rq  r>   rF   versionr@   r  r   r   getsourcelinesstarting_line_numberr  rO   r  r  rw  	enumerater.  r  r(   r   r   textwrapdedent	getsourceresearch	MULTILINEstart_unsafe_update_srcr   r  r  hashr/   r  r  rK   	arg_namesr   r   r  r  r   __globals__)r1   rq  r  r   r   r  rK   r   r  r   r   dnsdns_oar2   ps                  r4   r'   JITFunction.__init__W  s   1B-Ki)Goq&mm **2.!2.L+$+$:$:2$>q$A!
.!$..";";"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>>-- 7Hs   'JJ*Jc                t   U R                   c  [        U R                  U R                  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(   r-   r2   )r  r   rO   r  r2   ri   parser;   rR   r  dictsortedr/   r  )r1   dependencies_finders     r4   rQ   JITFunction.cache_key  s     99"4$--QUQaQagkgogo"p%%djjl3+//#d6O6O2PPDI$(0C0T0T0Z0Z0\)]$^D!yyr6   c               \    U R                   " [        [        R                  U5      USS.UD6$ )NTr  )r  r1  
MockTensor
wrap_dtype)r1   r  ru   r  s       r4   r  JITFunction.warmup  s(    xxZ5J5JD1QT$\U[\\r6   c           	     v   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                  R                  :w  a(  [        SUS    SU R                  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   )rr  r  r   r(   zSpecialization data is for z but trying to preload for r  r  r  r  r  r  r   )r  rr  r  r  triton.languager   r   r  r  loadsrq  rO   rN   r1  r   r/  r   is_dtyper  r  r%   r    r  )r1   ro  rr  r  r  tlrg  deserialized_objr  r  r   rn   r  r  r  r  r  r2   r  r  s                       r4   preloadJITFunction.preload  s   1$113::&9:F#tww'7'77-.>v.F-GGbcgcjcjcscsbtuw w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//F5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_   r  r2   r%   Moduler   bodyFunctionDef)r1   trees     r4   r  JITFunction.parse  se    yy"$

++++499~"""$))A,8888r6   c                    [        S5      e)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rN   )r1   ru   r  s      r4   __call__JITFunction.__call__  s    WXXr6   c                V   > US:X  a  [        SU S35      e[        [        U ]  X5        g )Nr2   zCannot set attribute 'zX' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)AttributeErrorr&   rL   __setattr__)r1   r(   rn   r3   s      r4   r  JITFunction.__setattr__  s:    5= #9$ @, "- . . 	k4,T9r6   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.
Nr2   )r  r&   r  )r1   new_srcr3   s     r4   r  JITFunction._unsafe_update_src  s    
 	E7+r6   c                P    SU R                    SU R                  R                   S3$ )NzJITFunction(:r$  )rF   rq  rO   r:   s    r4   __repr__JITFunction.__repr__  s&    dkk]!DGG,<,<+=Q??r6   )r  r  r   r  r>   rO   r  r  r  rr  r  r  r  r   r   rq  r  r  r  rF   rK   rw  r  r  r  r/   r  )NNNNNNN)rO   r>   r   r   ru  rv  r{  r  r  r  r   r'   r   rQ   r  r  r  r  r  r  r  r   r   r   s   @r4   rL   rL     s}    J M4
l(+ENF mq;?;(z  ]@Y:,@ @r6   rL   c                    g r8   r   )rq  s    r4   jitr    s    r6   r  r   r  r   r   r  rK   c                    g r8   r   r  s          r4   r  r    s     r6   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                  " SS5      S:X  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$ )NTRITON_INTERPRETr  r  r   )InterpretedFunction)r  r   r   r  rK   r   r  )r~  r  getenvinterpreterr  rL   )	rq  r  r  r   r   r  rK   r   r  s	     r4   	decoratorjit.<locals>.decorator  ss    |||99'-48&r7N_Fdlq08tUdf f "3/M! /	 	r6   rq  r   r   zJITFunction[T]r   )	rq  r  r   r  r   r   r  rK   r  s	    ``````` r4   r  r    s&    : & 
~} r6   c                  N    \ rS rSrSr\S 5       rS r\S 5       r\S 5       r	Sr
g)	r  i3  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)r3   rO   r>   r  )rv   s    r4   r  MockTensor.wrap_dtype9  s.    ==!!W,71Jc?"
r6   c                    Xl         g r8   r   )r1   r   s     r4   r'   MockTensor.__init__?  s    
r6   c                     gNr   r   r   r6   r4   data_ptrMockTensor.data_ptrB      r6   c                     gr  r   r   r6   r4   	ptr_rangeMockTensor.ptr_rangeF  r  r6   r  N)rO   r>   r   r   r   staticmethodr  r'   r	  r  r   r   r6   r4   r  r  3  sH    
  
    r6   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)TensorWrapperiK  c                    X l         Xl        UR                  U l        UR                  U l        U R                  R                  U l        g r8   )r   basedatarg  shape)r1   r  r   s      r4   r'   TensorWrapper.__init__M  s1    
	II	kkYY__
r6   c                6    U R                   R                  5       $ r8   )r  r	  r:   s    r4   r	  TensorWrapper.data_ptrT  s    yy!!##r6   c                4    U R                   R                  " U6 $ r8   )r  stride)r1   ru   s     r4   r  TensorWrapper.strideW  s    yy&&r6   c                <    SU R                    SU R                   S3$ )NzTensorWrapper[r`  r$  )r   r  r:   s    r4   __str__TensorWrapper.__str__Z  s    

|2dii[::r6   c                6    U R                   R                  5       $ r8   )r  element_sizer:   s    r4   r   TensorWrapper.element_size]  s    yy%%''r6   c                ^    [        U R                  R                  5       U R                  5      $ r8   )r  r  cpur   r:   s    r4   r#  TensorWrapper.cpu`  s    TYY]]_djj99r6   c                N    U R                   R                  UR                   5        g r8   )r  copy_)r1   others     r4   r&  TensorWrapper.copy_c  s    		

#r6   c                ^    [        U R                  R                  5       U R                  5      $ r8   )r  r  cloner   r:   s    r4   r*  TensorWrapper.clonef  s    TYY__.

;;r6   c                `    [        U R                  R                  U5      U R                  5      $ r8   )r  r  tor   )r1   rg  s     r4   r-  TensorWrapper.toi  s     TYY\\&14::>>r6   c                `    [        U R                  R                  U5      U R                  5      $ r8   )r  r  	new_emptyr   )r1   sizess     r4   r0  TensorWrapper.new_emptyl  s"    TYY007DDr6   )r  r  rg  r   r  Nr   rR   )rO   r>   r   r   r'   r	  r  r  r   r#  r&  r*  r-  r0  r   r   r6   r4   r  r  K  s5    %$';(:$<?Er6   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 r   )r%   r  r  r   r   r   r]   )r   r   s     r4   reinterpretr5  p  sm    &-((KK%%%;; !e44		$	$V++/V~Q?@@r6   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%   rL   rq  __code__co_filenamer@   r  r  striprC   )rq  base_fn	file_namelines
begin_lineidxlines          r4   get_jit_fn_file_liner@    s    G+..** +..

##//I..wzz:E u%	::<""6**J  	 &   r6   r3  )Fr  )r   Optional[Callable]r  rA  r   Optional[Iterable[int]]r   rB  r  Optional[bool]rK   rC  r   zCallable[[T], JITFunction[T]]r8   )rq  zOptional[T]r   rA  r  rA  r   rB  r   rB  r  rC  rK   rC  r   z4Union[JITFunction[T], Callable[[T], JITFunction[T]]])9
__future__r   r   r_   r)   r@   r   r  r  r  collectionsr   	functoolsr   typingr   r   r	   r
   r   r   r   r   r   r   runtime.driverr   typesr   _utilsr   r   rO   r   rD   r   NodeVisitorr   r   r   r   r   r   r  r  r"  r;  r   r    r  r  rL   r  r  r  r5  r@  r   r6   r4   <module>rL     s   , 
    	 	  # % d d d #  5.3~../CL~! ~!L-> ->` 	 /dZ
	Zgaj 	Z	0*f
D)  :	
 Y ) y 7 : z v  v v D  U!" U#$ - 2 
(//1	2A$%q! 
3a@/!$ a@R	 
 
 
 #*.15>B #
 
 (	

 /
 %<
 
 
 #
 

 4 #*.15>B #44 	4
 (4 /4 %<4 4 4 :4x 0"E "EJA!r6   