
    Ϧi              	         S SK JrJr  S SKrS SKrS SKrS SKrS SKrS SKrS SK	r	S SK
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	5      *  r\" S
5      r  " S S\RB                  5      r"SFS jr# " S S5      r$S r%0 r&SGS 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.S/_S0S1_S2S3_S4S5S6S7S8S9.Er+\," \+R[                  5       5       H  r.\.\+\.'   M
      " S: S;\(\    5      r/\SHS< j5       r0\SSSSSSSS=.             SIS> jj5       r0 SJSSSSSSSS=.               SKS? jjjr0 " S@ SA5      r1 " SB SC5      r2SD r3SE r4g)L    )annotationsdivisionN)defaultdict)cached_property)
CallableGenericIterableOptionalTypeVarUnionoverloadDictAnyTuple   )driver)
ModuleTypez.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       Q/var/www/html/ai-image-ml/venv/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r/   s    r2   retDependenciesFinder.retH   s    {{$$&&r4   c                    [         R                  " UR                  5      (       a  g[        USS5      nUR	                  [
        5      $ )NT
__module__ )inspect	isbuiltinfuncr"   
startswithTRITON_MODULE)r/   noder@   modules       r2   _is_triton_builtin%DependenciesFinder._is_triton_builtinL   s9    TYY''|R0  //r4   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)   )r/   r@   kvar_name_v1v2func_keys           r2   _update_hashDependenciesFinder._update_hashR   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 )r4   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+   getr.   r   r#   rJ   r"   r,   r-   rW   )r/   rC   vals      r2   
visit_NameDependenciesFinder.visit_Named   s    >SYY&77N77d&&&lltww-
 O 77 IZ/ #344WSJ^`e=f=fGG4#A#AABE||ATD!!477Bt||,<"=>#
r4   c                b    UR                    Vs/ s H  o R                  U5      PM     sn$ s  snf r6   )eltsvisit)r/   rC   elts      r2   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$ )NrM   r=   )	rg   valuer#   r]   	Attributer"   rB   attrrW   )r/   rC   lhsr9   s       r2   visit_Attribute"DependenciesFinder.visit_Attribute   s    jj$cmm,,**SYY'C cmm,,;73
B7=Hc99%#
r4   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 r6   )argsargr`   generic_visit)r/   rC   rt   s      r2   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.   rg   )defaultsexprr/   s     r2   visit_defaults:DependenciesFinder.visit_arguments.<locals>.visit_defaults   sN    8::::26/$D'

4( % 38/%/s    A A 	A)
	itertoolschainposonlyargsrs   vararg
kwonlyargsrg   kw_defaultskwargrz   )r/   rC   r|   rt   s   `   r2   visit_arguments"DependenciesFinder.visit_arguments   s    	8 ??4#3#3TYYQUQ\Q\bdfjfufuvCJJsO w 	t''(::!JJtzz"t}}%r4   c                    U R                  U5      n[        U[        5      (       a  U =R                  [	        U5      -  sl        g U R                  R                  U5        g r6   )rg   r#   r   r`   setadd)r/   rC   targets      r2   visitAssnTarget"DependenciesFinder.visitAssnTarget   sH     D!fd##F+  (r4   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   ru   r/   rC   s     r2   visit_AssignDependenciesFinder.visit_Assign   sG    t||!
 PQQT\\!_- 	4 r4   c                \    U R                  UR                  5        U R                  U5        g r6   r   r   ru   r   s     r2   visit_AnnAssign"DependenciesFinder.visit_AnnAssign   $    T[[) 	4 r4   c                \    U R                  UR                  5        U R                  U5        g r6   r   r   s     r2   	visit_ForDependenciesFinder.visit_For   r   r4   )r+   r*   r`   r&   r,   r-   r.   )returnNone)rM   r<   __qualname____firstlineno____doc__r%   propertyr9   rE   rW   rc   ri   rp   rv   r   r   r   r   r   __static_attributes____classcell__r1   s   @r2   r   r      s_    	"0H ' '09$<6
!
&@)!!! !r4   r   c                    [        U [        5      (       a  U R                  $ [        U [        5      (       a  U $ [	        U 5      $ r6   )r#   r[   rM   rP   repr)tys    r2   _normalize_tyr      s4    "d{{	B			8Or4   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 r6   )num_paramdo_not_specializedo_not_specialize_on_alignment)r/   r   paramr   r   s        r2   r%   KernelParam.__init__   s    !2.L+r4   c                .    U R                   R                  $ r6   )r   r&   r8   s    r2   r&   KernelParam.name   s    {{r4   c                    U R                   R                  (       a2  U R                   R                  [        R                  R                  :X  a  g[        U R                   R                  5      $ )Nr=   )r   
annotationr>   	Parameteremptyr   r8   s    r2   r   KernelParam.annotation   sD    {{%%)?)?7CTCTCZCZ)ZT[[3344r4   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   )r/   r   ty1ty2widths        r2   annotation_typeKernelParam.annotation_type   s\    __
5HCs3c#h>?@Eu*ug& 6 r4   c                     SU R                   ;   $ )N	constexpr)r   r8   s    r2   is_constexprKernelParam.is_constexpr  s    doo--r4   c                P    SU R                   ;   =(       a    U R                  (       + $ )Nconst)r   r   r8   s    r2   is_constKernelParam.is_const	  s    $//)C$2C2C.CCr4   c                .    U R                   R                  $ r6   )r   defaultr8   s    r2   r   KernelParam.default  s    {{"""r4   c                d    U R                   R                  [        R                  R                  :g  $ r6   )r   r   r>   r   r   r8   s    r2   has_defaultKernelParam.has_default  s#    {{""g&7&7&=&===r4   )r   r   r   r   N)r   r   r   zinspect.Parameterr   r   r   r   )rM   r<   r   r   r   r%   r   r&   r   r   r   r   r   r   r   r    r4   r2   r   r      s    LM15M     5 5
   . . D D # # > >r4   r   c                    U(       a)  [        U S5      (       a  U R                  5       S-  S:X  a  g[        U [        5      (       a  U(       a
  U S-  S:X  a  gU S:X  a  gg)Ndata_ptr   r   Dr   1N)hasattrr   r#   r   )valigns     r2   compute_spec_keyr     sP    J''QZZ\B->!-C	As		a"fk!Vr4   c                   U c  g[        U [        5      (       a  g[        U [        5      (       a  SU ::  a  U S::  a  gSU ::  a  U S::  a  gg	[        U [        5      (       a  g
[	        U S5      (       a  gU 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$ )Nnonei1   i32                u64i64fp32tma_desc_cpu_ptr	nvTmaDescr   *k*r   .)r#   r   r   r   r   dtype	dtype2strra   type_canonicalisation_dictrP   split)rt   r   dskress       r2   mangle_typer   &  s    
{	C			C		s?si/c\cY.	C			(	)	) yy(#mmC&;q64s.HSQRVIZIZ[^I_`bIc.ddC IcN
r4   c                  *    \ rS rSr% S\S'   SS jrSrg)KernelInterfaceiA  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   )rs   kwargsr   r/   s     r2   <lambda>-KernelInterface.__getitem__.<locals>.<lambda>J  s    txx$T%'YRX'Yr4   r   )r/   r   s   ``r2   __getitem__KernelInterface.__getitem__D  s     ZYr4   r   N)r   r   )rM   r<   r   r   __annotations__r  r   r   r4   r2   r   r   A  s    	
FZr4   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X#R                  5       UR                  WS.nUR                  U5      n	U	$ s  snnf )Nr   r   )r&   	signature	constantsattrsoptionskey)itemsr1   rM   rP   jsonto_dict__dict__dumps)
r&   r  r  r  r	  r
  rl   r  objserialized_objs
             r2   serialize_specialization_datar  N  sz    enetetevwevWaWZOO$<$<$Gc%jURevIw9}}C ZZ_N xs   3B c                6   [        U R                  5      [        U5      :X  d   e/ n/ n/ n/ n/ n/ n[        U R                  R                  5       U5       GHe  u  u  pnU
R                  [
        R                  R                  L a)  UR                  U	5        UR                  SU	 SU	 35        O-UR                  U	 SU	 35        UR                  SU	 SU	 35        UR                  (       a  UR                  U	5        M  UR                  U	5        UR                  (       d:  UR                  (       d  UR                  SU	-  5        OUR                  SU	-  5        UR                  (       a!  UR                  SUR                  -  5        GM8  UR                  SU	< SUR                  (       a  S	OS
< S35        GMh     SR                  Xx-    Vs/ s H  oS-   PM	     sn5      nSR                  U Vs/ s H  oS-   PM	     sn5      nSR                  U Vs/ s H  oS-   PM	     sn5      nUR                  S5        SR                  U5      nSR                  U5      nSU< SU< SU< SU< SU< S3nU R                  R                  5        V	Vs0 s H?  u  n	nUR                  [
        R                  R                  Ld  M/  SU	 3UR                  _MA     nn	n[        US'   UR                   US'   [#        UU5        US   $ s  snf s  snf 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': z	=default_z compute_spec_key(%s, align=True)z!compute_spec_key(%s, align=False)z"%s"zmangle_type(, TrueFalse)r=   z**excess_kwargszdef dynamic_func(z):
    return {z}, (z), (z), excess_kwargsdefault_r   r   dynamic_func)r   
parameterszipr  r   r>   r   r   appendr   r   r   r   r   joinr   r   exec)sigkparamsbackend	func_argsdict_entriesconstexpr_valsnon_constexpr_valssignature_typesspecialisationsr&   spkpxrO   args_strdict_str	func_bodyr   func_namespaces                      r2   create_function_from_signaturer0  Y  s    s~~#g,... ILNOO 4 4 6@$R::**000T"!D6TF 34vYtf56!D6TF 34??!!$'%%d+''88#**+MPT+TU#**+NQU+UV!!&&v0B0B'BC&&PRP[P[fahFh'ij' A* ?+LM+LaT+LMNIWW?1$h?@N4F!G4Fqd(4F!GH&' yy#Hyy&H(I~7IKI >>//11KD%== 1 1 7 77 	)(4&5==(1   %0N=!)0)A)AN%& 	N# .))5 N?!Gs   L4LL.LL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 5       r\SS j5       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S rSrU =r$ )rJ   i  Nc                0   [        U S5      (       a  U R                  $ [        U [        5      (       a  g[        U [        5      (       a  SU ::  a  U S::  a  gSU ::  a  U S::  a  gg	[        U [
        5      (       a  g
U c  g [        S[        U 5       SU  35      e)Nr   r   r   r   r   r   r   r   r   r   zUnsupported type z for )r   r   r#   r   r   r   r   r[   rt   s    r2   _key_ofJITFunction._key_of  s    3  99T""S!!33)#3##"2U##[/S	{%uEFFr4   c                    U c  g[        U [        5      (       a  U $ [        U 5      R                  S5      S   n[        U   nU(       a  SOSnX2-   $ )N*i8r   r   r   r   )r#   rP   r   r   )r
  r   	dtype_str	const_strs       r2   _type_ofJITFunction._type_of  sQ     ;S!!JHNN3'+	.y9	$D#	$$r4   c                D    [        [        U R                  U5      5      nU$ r6   )dictr  
constexprs)r/   constexpr_keyr  s      r2   _make_constantsJITFunction._make_constants  s    T__m<=	r4   c	                   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 S
3n " S S5      n[        XXFS   XQ5      nUUU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=](r  c                      \ rS rSrS rSrg)/JITFunction._call_hook.<locals>.JitFunctionInfoi  c                (    Xl         X l        X0l        g r6   )rD   r&   jit_function)r/   rD   r&   ri  s       r2   r%   8JITFunction._call_hook.<locals>.JitFunctionInfo.__init__  s    $ 	$0!r4   )ri  rD   r&   N)rM   r<   r   r   r%   r   r   r4   r2   JitFunctionInforg    s    r4   rk  r   )r  devicer  	num_warpsnum_ctas
num_stagesenable_fp_fusionextern_libsconfigsspecialization_data	is_warmupr
  )r
  r   fncompileis_manual_warmupalready_compiled)rJ   
cache_hookcompiled_hookru  rM   r<   r  r  paramsr&   rm  rn  ro  rp  r  rq  )r/   r
  r  rl  r  r	  rr  rt  beforehookr&   rD   r   r   	arg_reprsr   rk  rs  r   s                      r2   
_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  dm  cn  no  p	 	 <DY`aXbdkq #" **((!,, ' 8 8"..#6"
 vT2C*6*&"
 	
7 `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  )r/   r}  s     r2   add_pre_run_hookJITFunction.add_pre_run_hook  s&    
 ~~~!!$'r4   c                n   SSK JnJnJnJn  X l        X0l        X@l        XPl        [        U R                  U R                  U5      U l        [        U R                  5       VVs/ s H  u  pgUR                  (       d  M  UPM     snnU l        [        U R                  5       VVs/ s H  u  pgUR                  (       a  M  UPM     snnU l        [        U R                  5       VVs/ s H-  u  pgUR                  (       a  M  UR                  (       a  M+  UPM/     snnU l        gs  snnf s  snnf s  snnf )z!
Precompute as much as possible.
r   )CompiledKernelrv  	ASTSourcemake_backendN)compilerr  rv  r  r  r0  r  r{  binder	enumerater   constexpr_indicesnon_constexpr_indicesr   specialised_indices)r/   r"  r  rv  r  r  r   ps           r2   create_binderJITFunction.create_binder#  s     	PO,"(4T^^T[[RYZ2;DKK2H![2HANN!2H![6?6L%c6LFQTUTbTba6L%c"%dkk2$
2&11;N;NAYZYgYgA2$
  "\%c$
s*   #D%>D%%D+ D+'D1D1D1c               	   UR                  SS5      =(       d#    [        R                  R                  SS5      S:H  US'   SSKJn  [
        R                  R                  5       n[
        R                  R                  U5      n[
        R                  R                  5       nU" U5      n	U R                   H  n
U
" U0 UD6  M     U R                  c  U R                  U	5        U R                  " U0 UD6u  ppnSR                  U5      [        X45      -   nU R                  U   R                  US 5      nUGc  U	R!                  U5      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  [%        SU-  5      e   ['        UR)                  5       5      nU R*                   Vs/ s H  nU R,                  U   R.                  PM     nnUS [1        U5       n[3        UU5       VVs0 s H  u  nnUUS:X  a  SOU_M     nnnU	R5                  U R,                  U5      4nUS   R7                  5       n[3        UU R,                  5       VVs0 s H9  u  nnUR8                  (       d  UR:                  U;   d  Ub  M,  UR.                  U_M;     nnnUR=                  5        H%  u  nn[?        U5      (       d  M  [A        SU S35      e   U RC                  UUUUUUUSS9(       a  g U RE                  U UUUS   5      nU RG                  UUUR"                  S9nUU R                  U   U'   U RC                  UUUUUUUSS9  [I        5       n U RJ                  R=                  5        H8  u  u  n!n"u  n#n$U$R                  U!U 5      =n%U#:w  d  M&  [M        SU! SU# SU% 35      e   U(       d  Uc   e[?        U5      (       a  U" U5      n[1        U5      n&US   n'U&S:  a  US   OSn(U&S:  a  US   OSn)URN                  " X/UQ76 n*URP                  " U'U(U)UURR                  URT                  U*U RV                  RX                  U RV                  RZ                  /	UQ76   U$ s  snf s  snnf s  snnf )NdebugFTRITON_DEBUG0r   r   )r  r=   device_typez=device_type option is deprecated; current target will be usedrl  z8device option is deprecated; current device will be usedstreamz8stream option is deprecated; current stream will be usedz2Keyword argument %s was specified but unrecognisedr   rY  r   zCallable constexpr at index z is not supportedT)r|  )r   r	  rH   z1 has changed since we compiled this kernel, from z to r   ).ra   osenvironr  r  r   activeget_current_deviceget_current_streamget_current_targetr  r  r  r  rP   cacheparse_optionsr  KeyErrortuplevaluesr  r{  r&   r   r  get_attrs_descriptorget_constantsr   r   r  r  r   r  r  rv  objectr-   rL   launch_metadatar   functionpacked_metadatar  launch_enter_hooklaunch_exit_hook)+r/   r   r   rs   r   r  rl  r  r   r"  r}  
bound_argssig_and_specr%  r&  excess_kwargsr
  kernelr	  rQ   
bound_valsr   sigkeyssigvalsr   r  rr  constant_paramsr  r  rt   r0   not_presentr&   rS   rb   globals_dictnewVal	grid_sizegrid_0grid_1grid_2r  s+                                              r2   r   JITFunction.run3  s    **We4b

~WZ8[_b8bw 	,11311&9113v& &&D$!&! ' ;;w'VZVaVacgVrkqVrS
.m ggl#c>*I&JJF#''T2>++F3G !.o0oo.6)e+ee)6)e+ee)"G,,,"#WZ[#[\\ # z0023J 594N4NO4Nqt{{1~**4NGO"=CL1GJMgW^J_`J_AqF{U:J_I`33DKKLOG%aj668O "*dkk::FQ>>aee&>1 	:  
 $//+3C==#&B1#EV$WXX , sIvy'7TZcgh..y)WQZHC\\(( " F
 '-DJJvs#OOCFIwQW`eOf 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WDVWOJJvvvvvH^H^`o**<<d>Q>Q>b>byewyg P`s   ?#Q1Q6%+Q<Q<c	           	     0  ^^ U(       a  UO/ nU(       a  UO/ nTU l         TR                  U l        X l        [        R
                  " T5      U l        X0l        X@l        [        R                  " T5      S   U l	        UU4S jU l
        Xl        S U l        / 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,                  " T5      5      U l        U R.                  [0        R2                  " SU R.                  [0        R4                  5      R7                  5       S  U l        [9        [:        5      U l        S U l        0 U l         S U l!        X`l"        U R                   Vs/ s H  oR"                  PM     snU l#        U R                   Vs/ s H!  oRH                  (       d  M  URJ                  PM#     snU l&        / U l'        TRP                  U l(        TRR                  U l)        TRT                  U l*        TR                  U l        g s  snf s  snf )Nr   c                2   > Tc  TR                   $ T" U 5      $ r6   )rM   )rS   ru  r   s    r2   r   &JITFunction.__init__.<locals>.<lambda>  s    T\bkkFtAwFr4   z^def\s+\w+\s*\()+ru  r<   rD   versionr>   r  r   r   getsourcelinesstarting_line_numberr   r  r  r{  r  r  r  r&   r  r   textwrapdedent	getsourcer0   research	MULTILINEstartr   r_  r  hashr-   r  rI   	arg_namesr   r   r`  r  r   rM   __globals__)r/   ru  r  r   r   r  rI   r   r  r   r   dnsdns_oar  s    `     `      r2   r%   JITFunction.__init__  s   1B-Ki)Goq&mm **2.!2.L+$+$:$:2$>q$A!F	.!$..";";"B"B"DEHA(KEJJ:K,KC8hEJJJh<hFKK{1SAB F ??7#4#4R#8988BII&8$((BLLQWWYZ[ &
	 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+   r0   )r  r   rM   r  r0   rg   parser9   rP   r  r_  sortedr-   r  )r/   dependencies_finders     r2   rO   JITFunction.cache_key  s     99"4$--QUQaQagkgogo"p%%djjl3+//#d6O6O2PPDI$(0C0T0T0Z0Z0\)]$^D!yyr4   c               \    U R                   " [        [        R                  U5      USS.UD6$ )NTr   )r   map
MockTensor
wrap_dtype)r/   r   rs   r   s       r2   r   JITFunction.warmup  s(    xxZ5J5JD1QT$\U[\\r4   c           	     $   SSK JnJn  SSK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   R                  5        V	V
s0 s H8  u  pXR                   R#                  U
5      (       a  UR!                  U
5      OU
_M:     nn	n
[%        US	   R                  5       5      nU" XXR'                  US
   5      5      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XR.                  U   U	'   U$ s  sn
n	f s  sn
n	f )Nr   )rv  r  r   )AttrsDescriptorr&   zSpecialization data is for z but trying to preload for r  r  r  r	  r
  )r  rv  r  triton.backends.compilerr  r  triton.languagelanguager   r  r  loadsru  rM   rL   r  r   is_dtyper_  	from_dictr#   r   r  r  )r/   rs  rv  r  r  r  tlrl  deserialized_objr
  rl   r  r  r0   r	  r  s                   r2   preloadJITFunction.preload  s   1<$113::&9:F#tww'7'77-.>v.F-GGbcgcjcjcscsbtuw w /{;AAC
C
 HH$5$5e$<$<%%GC 	 
 )+6<<>?	4M4MN^_fNg4hi /y9??A
A
 E4!8!8ueCA 	 
 u%dG,"(

63

s   #?F2/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   )r]   r  r0   r#   Moduler   bodyFunctionDef)r/   trees     r2   r  JITFunction.parse  se    yy"$

++++499~"""$))A,8888r4   c                    [        S5      e)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rL   )r/   rs   r   s      r2   __call__JITFunction.__call__  s    WXXr4   c                H   > [         [        U ]  X5        US:X  a  S U l        g g )Nr0   )r$   rJ   __setattr__r  )r/   r&   rl   r1   s      r2   r  JITFunction.__setattr__   s'    k4,T9 5=DI r4   c                P    SU R                    SU R                  R                   S3$ )NzJITFunction(:r  )rD   ru  rM   r8   s    r2   __repr__JITFunction.__repr__  s&    dkk]!DGG,<,<+=Q??r4   )r  r  r   r  r<   rM   r  r  r  rv  r  r`  r   r   ru  r  r  r  r  rD   rI   r  r{  r  r   r  r  r0   r  r-   r  F)NNNNNNN)rM   r<   r   r   ry  rz  staticmethodrV  r\  rb  r  r  r  r   r%   r   rO   r   r  r  r  r  r  r   r   r   s   @r2   rJ   rJ     s    J MG G& 
% 
%3
j(
 \| mq;?:(x  ]8Y@ @r4   rJ   c                    g r6   r   )ru  s    r2   jitr    s    r4   r  r   r  r   r   r  rI   c                    g r6   r   r  s          r2   r  r    s     r4   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  rI   r   r  )r  r  getenvinterpreterr  rJ   )	ru  r  r  r   r   r  rI   r   r  s	     r2   	decoratorjit.<locals>.decorator@  ss    |||99'-48&r7N_Fdlq08tUdf f "3/M! /	 	r4   ru  r   r   zJITFunction[T]r   )	ru  r  r   r  r   r   r  rI   r  s	    ``````` r2   r  r  #  s&    : & 
~} r4   c                  N    \ rS rSrSr\S 5       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)r1   rM   r<   r  rU  s    r2   r  MockTensor.wrap_dtypee  s.    ==!!W,71Jc?"
r4   c                    Xl         g r6   r   )r/   r   s     r2   r%   MockTensor.__init__k  s    
r4   c                     gNr   r   r   r4   r2   r   MockTensor.data_ptrn      r4   c                     gr  r   r   r4   r2   	ptr_rangeMockTensor.ptr_ranger  r  r4   r	  N)rM   r<   r   r   r   r  r  r%   r   r  r   r   r4   r2   r  r  _  sH    
  
    r4   r  c                  N    \ 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g)TensorWrapperiw  c                    X l         Xl        UR                  U l        UR                  U l        U R                  R                  U l        g r6   )r   basedatarl  shape)r/   r  r   s      r2   r%   TensorWrapper.__init__y  s1    
	II	kkYY__
r4   c                6    U R                   R                  5       $ r6   )r  r   r8   s    r2   r   TensorWrapper.data_ptr  s    yy!!##r4   c                8    U R                   R                  U5      $ r6   )r  stride)r/   r   s     r2   r  TensorWrapper.stride  s    yy""r4   c                <    SU R                    SU R                   S3$ )NzTensorWrapper[re  r  )r   r  r8   s    r2   __str__TensorWrapper.__str__  s    

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

#r4   c                ^    [        U R                  R                  5       U R                  5      $ r6   )r  r  cloner   r8   s    r2   r,  TensorWrapper.clone  s    TYY__.

;;r4   c                `    [        U R                  R                  U5      U R                  5      $ r6   )r  r  tor   )r/   rl  s     r2   r/  TensorWrapper.to  s     TYY\\&14::>>r4   )r  r  rl  r   r  Nr   rP   )rM   r<   r   r   r%   r   r  r  r"  r%  r(  r,  r/  r   r   r4   r2   r  r  w  s/    %$#;(:$<?r4   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[   )tensorr   s     r2   reinterpretr4    sm    &-((KK%%%;; !e44		$	$V++/V~Q?@@r4   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#   rJ   ru  __code__co_filenamer>   r  r  striprA   )ru  base_fn	file_namelines
begin_lineidxlines          r2   get_jit_fn_file_liner?    s    G+..** +..

##//I..wzz:E u%	::<""6**J  	 &   r4   r1  r  r  )r   Optional[Callable]r  r@  r   Optional[Iterable[int]]r   rA  r  Optional[bool]rI   rB  r   zCallable[[T], JITFunction[T]]r6   )ru  zOptional[T]r   r@  r  r@  r   rA  r   rA  r  rB  rI   rB  r   z4Union[JITFunction[T], Callable[[T], JITFunction[T]]])5
__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   rM   r   rB   r   NodeVisitorr   r   r   r   r   r   r   r  r0  r   r   r  r   rJ   r  r  r  r4  r?  r   r4   r2   <module>rJ     s   , 
    	 	  # % d d d # .3~../CL~! ~!L-> ->`
 	6	Zgaj 	ZA*H
D)  :	
 Y ) y 7 : z v  v v D  U!" U#$ - 2 
(//1	2A$%q! 
3N@/!$ N@l
 
 
 
 #*.15>B #
 
 (	

 /
 %<
 
 
 #
 

 4 #*.15>B #44 	4
 (4 /4 %<4 4 4 :4x 0? ?DA!r4   