
    "h              	         d dl mZmZ d dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
mZ d dlmZ d dlmZmZmZmZmZmZmZmZmZmZ ddlmZ d dlmZ ed ed	        Z ed
      Z  G d dejB                        Z"dFdZ# G d d      Z$d Z%i Z&dGdZ' G d dee          Z(d Z)d Z*i dddddddddddd d!d d"dd#d$d%d$d&d'd(d)d*d+d,d-d.d/d0d1d2d3d4d5d6d7d8d9Z+ e,e+j[                               D ]  Z.e.e+e.<   	  G d: d;e(e          Z/edHd<       Z0edddddddd=	 	 	 	 	 	 	 	 	 	 	 	 	 dId>       Z0	 dJdddddddd=	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 dKd?Z0 G d@ dA      Z1 G dB dC      Z2dD Z3dE Z4y)L    )annotationsdivisionN)defaultdict)cached_property)
CallableGenericIterableOptionalTypeVarUnionoverloadDictAnyTuple   )driver)
ModuleTypez.runtime.jitTc                  v     e Zd ZdZd fdZed        Zd Zd Zd Z	d Z
d Zd	 Zd
 Zd Zd Zd Zd Z xZS )DependenciesFindera  
    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         |           || _        t        j                  |j                  d            | _        || _        h d| _        i | _	        d| _
        y )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       O/var/www/html/sandstorm/venv/lib/python3.12/site-packages/triton/runtime/jit.pyr$   zDependenciesFinder.__init__$   sU    	nnSZZ%89 *
&. TV*/'    c                6    | j                   j                         S N)r)   	hexdigestr.   s    r1   retzDependenciesFinder.retH   s    {{$$&&r2   c                    t        j                  |j                        ryt        |dd      }|j	                  t
              S )NT
__module__ )inspect	isbuiltinfuncr!   
startswithTRITON_MODULE)r.   noder=   modules       r1   _is_triton_builtinz%DependenciesFinder._is_triton_builtinL   s6    TYY'|R0  //r2   c                >   t        |t              r| j                  j                         |j                  j                         z  D ]_  }|\  }}| j                  |   \  }}|j                  |   \  }}||k7  s2t	        d| d| d| j
                   d|j                   d| d       | j                  j                  |j                         |j                  }|t        t        |dd            z  }| j                  j                  |j                  d	             y y )
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           r1   _update_hashzDependenciesFinder._update_hashR   s3   dK( **//1D4I4I4N4N4PP !--a0A--a0A8&*8*KtCSTXT]T]S^^qrvrr  rA  AX  Y[  X\  \S  T  !!(()>)>?~~HGD*e<==HKKxw78 )r2   c                $   t        |j                        t        j                  u r|j                  S |j                  | j
                  v ry | j                  j                  |j                  d       }|| j                  st        |      t        urot        |t              s_t        |dd      sR|j                  | j                  vr:|| j                  f| j                  |j                  t	        | j                        f<   | j                  |       |S )N__triton_builtin__F)typectxastStoreidlocal_namesr*   getr-   r   r"   rF   r!   r+   r,   rS   )r.   r@   vals      r1   
visit_NamezDependenciesFinder.visit_Named   s    >SYY&77N77d&&&lltww-
 O 77 IZ/ #34WSJ^`e=fGG4#A#AABEt||ATD!!477Bt||,<"=>#
r2   c                ^    |j                   D cg c]  }| j                  |       c}S c c}w r4   )eltsvisit)r.   r@   elts      r1   visit_TuplezDependenciesFinder.visit_Tuple   s$     ,0995C

3555s   *c                X   | j                  |j                        }t        |t        j                        r6| j                  |j                        }t        |t        j                        r6|t        |dd      t        k(  ry t        ||j                        }| j                  |       |S )NrI   r:   )	ra   valuer"   rX   	Attributer!   r?   attrrS   )r.   r@   lhsr7   s       r1   visit_Attributez"DependenciesFinder.visit_Attribute   s    jj$cmm,**SYY'C cmm,;73
B7=Hc499%#
r2   c                    |j                   j                   D ch c]  }|j                   c}| _        | j                  |       y c c}w r4   )argsargr[   generic_visit)r.   r@   rl   s      r1   visit_FunctionDefz$DependenciesFinder.visit_FunctionDef   s4    /3yy~~>CGG>4  ?s   Ac                p     fd}t        j                  |j                  |j                  |j                  r|j                  gng |j
                        D ]  } j                  |         ||j                         |j                   j                  |j                          ||j                         y )Nc                    	 j                   rJ d_         | D ]  }|j                  |        	 d_         y # d_         w xY w)NTF)r-   ra   )defaultsexprr.   s     r1   visit_defaultsz:DependenciesFinder.visit_arguments.<locals>.visit_defaults   sS    8::::26/$ )D'

4() 38/%/s   < < 	A)
	itertoolschainposonlyargsrk   vararg
kwonlyargsra   kw_defaultskwargrq   )r.   r@   rs   rl   s   `   r1   visit_argumentsz"DependenciesFinder.visit_arguments   s    	8 ??4#3#3TYYQUQ\Q\bdfjfufuv 	CJJsO	 	t''(::!JJtzz"t}}%r2   c                    | j                  |      }t        |t              r| xj                  t	        |      z  c_        y | j                  j                  |       y r4   )ra   r"   r   r[   setadd)r.   r@   targets      r1   visitAssnTargetz"DependenciesFinder.visitAssnTarget   sE     D!fd#F+  (r2   c                    t        |j                        dk7  rt        d      | j                  |j                  d          | j	                  |       y )N   z2Simultaneous multiple assignment is not supported.r   )r   targets	TypeErrorr   rm   r.   r@   s     r1   visit_AssignzDependenciesFinder.visit_Assign   sG    t||!
 PQQT\\!_- 	4 r2   c                \    | j                  |j                         | j                  |       y r4   r   r   rm   r   s     r1   visit_AnnAssignz"DependenciesFinder.visit_AnnAssign   $    T[[) 	4 r2   c                \    | j                  |j                         | j                  |       y r4   r   r   s     r1   	visit_ForzDependenciesFinder.visit_For   r   r2   )returnNone)rI   r9   __qualname____doc__r$   propertyr7   rB   rS   r^   rc   ri   rn   r{   r   r   r   r   __classcell__r0   s   @r1   r   r      sZ    	"0H ' '09$<6
!
&@)!!!r2   r   c                t    t        | t              r| j                  S t        | t              r| S t	        |       S r4   )r"   rV   rI   rL   repr)tys    r1   _normalize_tyr      s.    "d{{	B		8Or2   c                      e Zd ZdZ	 	 ddZed        Zed        Zed        Zed        Z	ed        Z
ed        Zed	        Zy
)KernelParamzBRepresents a parameter (name plus metadata) to a @jit'ed function.c                <    || _         || _        || _        || _        y r4   )num_paramdo_not_specializedo_not_specialize_on_alignment)r.   r   paramr   r   s        r1   r$   zKernelParam.__init__   s     !2.L+r2   c                .    | j                   j                  S r4   )r   r%   r6   s    r1   r%   zKernelParam.name   s    {{r2   c                    | j                   j                  r1| j                   j                  t        j                  j                  k(  ryt        | j                   j                        S )Nr:   )r   
annotationr;   	Parameteremptyr   r6   s    r1   r   zKernelParam.annotation   sD    {{%%)?)?7CTCTCZCZ)ZT[[3344r2   c                    | j                   }dD ]4  \  }}||j                  |      t        |      z   d  }|s)||v s.| | c S  |dk(  ryy)N))uintu)r   iboolu1r:   )r   findr   )r.   r   ty1ty2widths        r1   annotation_typezKernelParam.annotation_type   se    __
5 	'HCzs3c#h>?@E
*ug&	' r2   c                    d| j                   v S )N	constexpr)r   r6   s    r1   is_constexprzKernelParam.is_constexpr  s    doo--r2   c                <    d| j                   v xr | j                   S )Nconst)r   r   r6   s    r1   is_constzKernelParam.is_const	  s    $//)C$2C2C.CCr2   c                .    | j                   j                  S r4   )r   defaultr6   s    r1   r   zKernelParam.default  s    {{"""r2   c                d    | j                   j                  t        j                  j                  k7  S r4   )r   r   r;   r   r   r6   s    r1   has_defaultzKernelParam.has_default  s#    {{""g&7&7&=&===r2   N)r   r   r   zinspect.Parameterr   r   r   r   )rI   r9   r   r   r$   r   r%   r   r   r   r   r   r   r    r2   r1   r   r      s    LM15M     5 5
   . . D D # # > >r2   r   c                    |r#t        | d      r| j                         dz  dk(  ryt        | t              r|r	| dz  dk(  ry| dk(  ryy)Ndata_ptr   r   Dr   1N)hasattrr   r"   r   )valigns     r1   compute_spec_keyr     sJ    J'QZZ\B->!-C	As	a"fk!Vr2   c                x   | yt        | t              ryt        | t              rd| k  r| dk  ryd| k  r| dk  ryy	t        | t              ry
t	        | d      ry| j
                  |f}t        j                  |d       }|:|d   rdndt        t        |d         j                  d      d      z   }|t        |<   |S )Nnonei1   i32                u64i64fp32tma_desc_cpu_ptr	nvTmaDescr   *k*r   .)r"   r   r   r   r   dtype	dtype2strr\   type_canonicalisation_dictrL   split)rl   r   dskress       r1   mangle_typer   &  s    
{	C		C	s?si/c\cY.	C		(	) yy(#mmC&;q64s.HSQRVIZIZ[^I_`bIc.ddC IcN
r2   c                       e Zd ZU ded<   ddZy)KernelInterfacer   runc                      fdS )z
        A JIT function is launched with: fn[grid](*args, **kwargs).
        Hence JITFunction.__getitem__ returns a callable proxy that
        memorizes the grid.
        c                 .     j                   | dd|S )NFgridwarmup)r   )rk   kwargsr   r.   s     r1   <lambda>z-KernelInterface.__getitem__.<locals>.<lambda>J  s    xtxx$T%'YRX'Y r2   r   )r.   r   s   ``r1   __getitem__zKernelInterface.__getitem__D  s     ZYr2   N)r   r   )rI   r9   r   __annotations__r   r   r2   r1   r   r   A  s    	
FZr2   r   c                   |j                         D ci c],  \  }}||j                  j                  dk(  rt        |      n|. }}}dd l}| |||j                         |j                  d}|j                  |      }	|	S c c}}w )Nr   r   )r%   	signature	constantsattrsoptionskey)itemsr0   rI   rL   jsonto_dict__dict__dumps)
r%   r   r   r   r   r   re   r   objserialized_objs
             r1   serialize_specialization_datar   N  s    enetetevwWaWZ\aEOO$<$<$Gc%jURwIw99u}}C ZZ_N xs   1B c                   t        | j                        t        |      k(  sJ g }g }g }g }g }g }t        | j                  j                         |      D ]G  \  \  }	}
}|
j                  t
        j                  j                  u r)|j                  |	       |j                  d|	 d|	        n-|j                  |	 d|	        |j                  d|	 d|	        |j                  r|j                  |	       |j                  |	       |j                  s5|j                  s|j                  d|	z         n|j                  d|	z         |j                  r |j                  d|j                  z         "|j                  d|	d|j                  rd	nd
d       J dj                  ||z   D cg c]  }|dz   	 c}      }dj                  |D cg c]  }|dz   	 c}      }dj                  |D cg c]  }|dz   	 c}      }|j                  d       dj                  |      }dj                  |      }d|d|d|d|d|d}| j                  j                         D 	ci c];  \  }	}|j                  t
        j                  j                  urd|	 |j                  = }}	}t        |d<   |j                   |d<   t#        ||       |d   S c c}w c c}w c c}w c c}}	w )a2  
    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xrK   args_strdict_str	func_bodyr   func_namespaces                      r1   create_function_from_signaturer  Y  s    s~~#g,... ILNOO 4 4 6@ k$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'k* ?_+LMaTMNIWW?1a$h?@N4F!Gq!d(!GH&' yy#Hyy&H(I~7IKI >>//1D%== 1 1 7 77 4&5==(N  %0N=!)0)A)AN%& 	N# .))5 N?!Gs   1KK5K *A K%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                       e Zd ZdZdZed        Zedd       Zd Zd Z	d Z
d Zd Z	 	 dd	Zed
        Zd Zd Zd Zd Z fdZd Z xZS )rF   Nc                    t        | d      r| j                  S t        | t              ryt        | t              rd| k  r| dk  ryd| k  r| dk  ryy	t        | t
              ry
| y t        dt        |        d|        )Nr   r   r   r   r   r   r   r   r   r   zUnsupported type z for )r   r   r"   r   r   r   r   rV   rl   s    r1   _key_ofzJITFunction._key_of  s    3 99T"S!33)#3##"2U#[/S	{%uEFFr2   c                    | yt        | t              r| S t        |       j                  d      d   }t        |   }|rdnd}||z   S )N*i8r   r   r   r   )r"   rL   r   r   )r   r   	dtype_str	const_strs       r1   _type_ofzJITFunction._type_of  sP     ;S!JHNN3'+	.y9	$D#	9$$r2   c                D    t        t        | j                  |            }|S r4   )dictr   
constexprs)r.   constexpr_keyr   s      r1   _make_constantszJITFunction._make_constants  s    T__m<=	r2   c	                   |rt         j                  nt         j                  }	|	y| j                  j                  }
| j                  j
                  }dj                  t        | j                  |d         D cg c]  \  }}|j                   d|  c}}      }|
 d|j                   d|j                   d|j                   d|j                   d	| d
} G d d      }t        |
|||d   ||      }||||j                  |j                  |j                  |j                  |j                  |||d} |	|| |||
|       d|i||d      S c c}}w )NFr   r   z: z[num_warps=z, num_ctas=z, num_stages=z, enable_fp_fusion=](r   c                      e Zd Zd Zy)/JITFunction._call_hook.<locals>.JitFunctionInfoc                .    || _         || _        || _        y r4   )rA   r%   jit_function)r.   rA   r%   rJ  s       r1   r$   z8JITFunction._call_hook.<locals>.JitFunctionInfo.__init__  s    $ 	$0!r2   N)rI   r9   r   r$   r   r2   r1   JitFunctionInforH    s    r2   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)rF   
cache_hookcompiled_hookrU  rI   r9   r  r   paramsr%   rM  rN  rO  rP  r   rQ  )r.   r   r   rL  r   r   rR  rT  beforehookr%   rA   r   r   	arg_reprsr   rK  rS  r   s                      r1   
_call_hookzJITFunction._call_hook  s    *0{%%[5N5N<ww##IIc$++WZ[\W]F^_%**Rt4_`	{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	 	 <D)YX_`aXbdkmpq #" **((!,, ' 8 8"..#6"
 vtT2C*6*&"
 	
7 `s   9E
c                T    t        |      sJ | j                  j                  |       y)z
        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     r1   add_pre_run_hookzJITFunction.add_pre_run_hook  s$    
 ~~!!$'r2   c                :   ddl m}m}m}m} || _        || _        || _        || _        t        | j                  | j                  |      | _        t        | j                        D cg c]  \  }}|j                  s| c}}| _        t        | j                        D cg c]  \  }}|j                  r| c}}| _        t        | j                        D cg c]!  \  }}|j                  r|j                  r |# c}}| _        yc c}}w c c}}w c c}}w )z1
        Precompute as much as possible.
        r   )CompiledKernelrV  	ASTSourcemake_backendN)compilerre  rV  rf  rg  r  r   r[  binder	enumerater   constexpr_indicesnon_constexpr_indicesr   specialised_indices)r.   r  re  rV  rf  rg  r   ps           r1   create_binderzJITFunction.create_binder#  s     	PO,"(4T^^T[[RYZ2;DKK2H![AANN!![6?6L%cFQTUTbTba%c"%dkk2$
1a1;N;NYZYgYgA$
  "\%c$
s*   'D<D!D6DD0D=Dc                  |j                  dd      xs# t        j                  j                  dd      dk(  |d<   ddlm} t
        j                  j                         }t
        j                  j                  |      }t
        j                  j                         } ||      }	| j                  D ]
  }
 |
|i |  | j                  | j                  |	        | j                  |i |\  }}}}}dj                  |      t        ||f      z   }| j                  |   j                  |d       }||	j!                  |      }d	|vsJ d
       d|vsJ d       d|vsJ d       |D ]  }||j"                  vst%        d|z         t'        |j)                               }| j*                  D cg c]  }| j,                  |   j.                   }}|d t1        |       }t3        ||      D ci c]  \  }}||dk(  rdn| }}}|	j5                  | j,                  |      f}|d   j7                         }t3        || j,                        D ci c].  \  }}|j8                  s|j:                  |v s||j.                  |0 }}}|j=                         D ]  \  }}t?        |      stA        d| d       | jC                  |||||||d      ry | jE                  | |||d         }| jG                  |||j"                        }|| j                  |   |<   | jC                  |||||||d       tI               } | jJ                  j=                         D ]6  \  \  }!}"\  }#}$|$j                  |!|       x}%|#k7  s$tM        d|! d|# d|%        |s|J t?        |      r ||      }t1        |      }&|d   }'|&dkD  r|d   nd}(|&dkD  r|d   nd}) |jN                  ||g| }* |jP                  |'|(|)||jR                  |jT                  |*| jV                  jX                  | jV                  jZ                  g	|  |S c c}w c c}}w c c}}w )NdebugFTRITON_DEBUG0r   r   )rg  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   r<  r   zCallable constexpr at index z is not supportedT)r\  )r   r   rD   z1 has changed since we compiled this kernel, from z to r   ).r\   osenvironrh  rg  r   activeget_current_deviceget_current_streamget_current_targetrb  ri  ro  r  rL   cacheparse_optionsr   KeyErrortuplevaluesrl  r[  r%   r   r   get_attrs_descriptorget_constantsr   r   r   ra  r   r_  rf  rV  objectr,   rH   launch_metadatar   functionpacked_metadatare  launch_enter_hooklaunch_exit_hook)+r.   r   r   rk   r   rg  rL  ru  r   r  r]  
bound_argssig_and_specr	  r
  excess_kwargsr   kernelr   rM   
bound_valsr   sigkeyssigvalsr   r   rR  constant_paramsrn  r   rl   r/   not_presentr%   rO   r]   globals_dictnewVal	grid_sizegrid_0grid_1grid_2r  s+                                              r1   r   zJITFunction.run3  s    **We4b

~WZ8[_b8bw 	,11311&9113v& && 	"D$!&!	" ;;w'VaVZVaVacgVrkqVrS
L.2Dm 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Oqt{{1~**OGO"=CL1GJMgW^J_`AqF{U:`I`33DKKLOG%aj668O "*dkk:Q>>aee&>19 	I 
 $//+ Y3C=#&B1#EV$WXXY sIvy'7TZcgh..y)WQZHC\\(( " F
 '-DJJvs#OOCFIwQW`eOf h.2.C.C.I.I.K 	q*IT1*\&**4==#E"&tf,]^a]bbfgmfnoq q	q
 ###~ J'D	I!WF )AT!W1F )AT!W1F 5f44T6WDVWOFJJvvvvvH^H^`o**<<d>Q>Q>b>byewyg P`s   / P;.Q 3Qc	           	        |r|ng }|r|ng }| _         j                  | _        || _        t	        j
                        | _        || _        || _        t	        j                        d   | _	        fd| _
        || _        d | _        g | _        t        | j
                  j                  j!                               D ]T  \  }	}
|	|v xs |
j"                  |v }|	|v xs |
j"                  |v }| j                  j%                  t'        |	|
||             V t)        j*                  t	        j,                              | _        | j.                  t1        j2                  d| j.                  t0        j4                        j7                         d  | _        t9        t:              | _        d | _        i | _         d | _!        || _"        | j                  D cg c]  }|j"                   c}| _#        | j                  D cg c]  }|jH                  s|jJ                   c}| _&        g | _'        jP                  | _(        jR                  | _)        jT                  | _*        j                  | _        y c c}w c c}w )Nr   c                0    j                   S  |       S r4   )rI   )rO   rU  r   s    r1   r   z&JITFunction.__init__.<locals>.<lambda>  s    T\bkk tAw r2   z^def\s+\w+\s*\()+rU  r9   rA   versionr;   r   r   r   getsourcelinesstarting_line_numberr   r  ri  r[  rj  r   r  r%   r  r   textwrapdedent	getsourcer/   research	MULTILINEstartr   rA  r|  hashr,   r  rE   	arg_namesr   r   rB  rb  r   rI   __globals__)r.   rU  r  r   r   rq  rE   r   r  r   r   dnsdns_oarn  s    `     `      r1   r$   zJITFunction.__init__  s   1B-Ki)Goq&mm **2.!2.L+$+$:$:2$>q$A!F	.!$..";";"B"B"DE 	CHAu((KEJJ:K,KC88hEJJJh<hFKK{1eS&AB	C ??7#4#4R#8988BII&8$((BLLQWWYZ[ &
	 TV   +/++6Q!&&6*.++HQ155H   zz>>-- 7Hs   I.:I3I3c                x   | j                   t        | j                  | j                  | j                        }|j                  | j                                |j                  t        | j                        z   | _         t        t        |j                  j                                     | _        | j                   S )N)r%   r*   r/   )r  r   rI   r  r/   ra   parser7   rL   r  rA  sortedr,   r   )r.   dependencies_finders     r1   rK   zJITFunction.cache_key  s     99"4$--QUQaQagkgogo"p%%djjl3+//#d6O6O2PPDI$(0C0T0T0Z0Z0\)]$^D!yyr2   c               \     | j                   t        t        j                  |      |dd|S )NTr   )r   map
MockTensor
wrap_dtype)r.   r   rk   r   s       r1   r   zJITFunction.warmup  s*    txxZ5J5JD1QT$\U[\\r2   c           	        ddl m}m} ddlm} dd l}dd lm} t        j                  j                         }|j                  |      }|d   | j                  j                  k7  r(t        d|d    d| j                  j                         |d   j                         D 	
ci c]4  \  }	}
|	|j                   j#                  |
      r|j!                  |
      n|
6 }}	}
t%        |d	   j                               } || |||j'                  |d
               }|d   j                         D 	
ci c]#  \  }	}
|	t)        |
t*              rt-        |
      n|
% }}	}
|d   }	 ||d |      }|| j.                  |   |	<   |S c c}
}	w c c}
}	w )Nr   )rV  rf  r   )AttrsDescriptorr%   zSpecialization data is for z but trying to preload for r   r   r   r   r   )rh  rV  rf  triton.backends.compilerr  r   triton.languagelanguager   rx  ry  loadsrU  rI   rH   r   r   is_dtyperA  	from_dictr"   r   r  r|  )r.   rS  rV  rf  r  r   tlrL  deserialized_objr   re   r   r   r/   r   r  s                   r1   preloadzJITFunction.preload  s   1<$113::&9:F#tww'7'77-.>v.F-GGbcgcjcjcscsbtuw w /{;AAC
U BHH$5$5e$<%%G
	 
 )+6<<>?	iO4M4MN^_fNg4hi /y9??A
U E4!8ueC
 
 u%dG,"(

63

s   "9E;-(Fc                   t        j                  | j                        }t        |t         j                        sJ t        |j                        dk(  sJ t        |j                  d   t         j                        sJ |S )Nr   r   )rX   r  r/   r"   Moduler   bodyFunctionDef)r.   trees     r1   r  zJITFunction.parse  s_    yy"$

+++499~"""$))A,888r2   c                    t        d      )Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rH   )r.   rk   r   s      r1   __call__zJITFunction.__call__  s    WXXr2   c                H    t         t        |   ||       |dk(  rd | _        y y )Nr/   )r#   rF   __setattr__r  )r.   r%   re   r0   s      r1   r  zJITFunction.__setattr__   s)    k4,T59 5=DI r2   c                P    d| j                    d| j                  j                   dS )NzJITFunction(:r   )rA   rU  rI   r6   s    r1   __repr__zJITFunction.__repr__  s&    dkk]!DGG,<,<+=Q??r2   F)NNNNNNN)rI   r9   r   rY  rZ  staticmethodr:  r?  rD  r_  rc  ro  r   r$   r   rK   r   r  r  r  r  r  r   r   s   @r1   rF   rF     s    J MG G& 
% 
%3
j(
 \| mq;?:(x  ]8Y@r2   rF   c                     y r4   r   )rU  s    r1   jitr    s    r2   r  r   r  r   r   rq  rE   c                     y r4   r   r  s          r1   r  r    s     r2   c               @    dfd}|  ||       S |S )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           
         t        |       sJ t        j                  dd      dk(  rddlm}  ||       S t        |       S )NTRITON_INTERPRETrs  r   r   )InterpretedFunction)r  r   r   rq  rE   r   r  )ra  rv  getenvinterpreterr  rF   )	rU  r  rq  r   r   r  rE   r   r  s	     r1   	decoratorzjit.<locals>.decorator@  sq    ||99'-48&r7N_Fdlq08tUdf f "3/M! /	 	r2   rU  r   r   zJITFunction[T]r   )	rU  r  r   r  r   r   rq  rE   r  s	    ``````` r1   r  r  #  s&    : & 
~} r2   c                  F    e Zd ZdZed        Zd Zed        Zed        Zy)r  zr
    Can be used in place of real tensors when calling:
        kernel.warmup(MockTensor(torch.float32), ...)
    c                l    | j                   j                  dk(  r| j                  dk(  rt        |       S | S )Nr   torch)r0   rI   r9   r  r9  s    r1   r  zMockTensor.wrap_dtypee  s.    ==!!W,71Jc?"
r2   c                    || _         y r4   )r   )r.   r   s     r1   r$   zMockTensor.__init__k  s	    
r2   c                      yNr   r   r   r2   r1   r   zMockTensor.data_ptrn      r2   c                      yr  r   r   r2   r1   	ptr_rangezMockTensor.ptr_ranger  r  r2   N)	rI   r9   r   r   r  r  r$   r   r  r   r2   r1   r  r  _  sH    
  
    r2   r  c                  D    e Zd Zd Zd Zd ZddZd Zd Zd Z	d Z
d	 Zy
)TensorWrapperc                    || _         || _        |j                  | _        |j                  | _        | j                  j                  | _        y r4   )r   basedatarL  shape)r.   r  r   s      r1   r$   zTensorWrapper.__init__y  s5    
	II	kkYY__
r2   c                6    | j                   j                         S r4   )r  r   r6   s    r1   r   zTensorWrapper.data_ptr  s    yy!!##r2   c                8    | j                   j                  |      S r4   )r  stride)r.   r   s     r1   r  zTensorWrapper.stride  s    yy""r2   c                <    d| j                    d| j                   dS )NzTensorWrapper[rF  r   )r   r  r6   s    r1   __str__zTensorWrapper.__str__  s    

|2dii[::r2   c                6    | j                   j                         S r4   )r  element_sizer6   s    r1   r  zTensorWrapper.element_size  s    yy%%''r2   c                ^    t        | j                  j                         | j                        S r4   )r  r  cpur   r6   s    r1   r  zTensorWrapper.cpu  s    TYY]]_djj99r2   c                N    | j                   j                  |j                          y r4   )r  copy_)r.   others     r1   r  zTensorWrapper.copy_  s    		

#r2   c                ^    t        | j                  j                         | j                        S r4   )r  r  cloner   r6   s    r1   r  zTensorWrapper.clone  s    TYY__.

;;r2   c                `    t        | j                  j                  |      | j                        S r4   )r  r  tor   )r.   rL  s     r1   r  zTensorWrapper.to  s     TYY\\&14::>>r2   Nr   rL   )rI   r9   r   r$   r   r  r  r  r  r  r  r  r   r2   r1   r  r  w  s/    %$#;(:$<?r2   r  c                    t        | t              r;|| j                  j                  k(  r| j                  S t        | j                  |      S t	        | d      rt        | |      S t        dt        |        d      )Nr   zCannot reinterpret a r   )r"   r  r  r   r   r   rV   )tensorr   s     r1   reinterpretr    sk    &-(FKK%%%;; !e44		$VU++/V~Q?@@r2   c                d   | }t        |t              s|j                  }t        |t              s|j                  j                  j                  }t        j                  |j                        \  }}t        |      D ].  \  }}|j                         j                  d      s&||z  } ||fS  ||fS )Nzdef )
r"   rF   rU  __code__co_filenamer;   r  rj  stripr>   )rU  base_fn	file_namelines
begin_lineidxlines          r1   get_jit_fn_file_liner    s    G+.** +.

##//I..wzz:E: u% 	T::<""6*#Jj  	 j  r2   r  r  r  )r   Optional[Callable]r  r  r   Optional[Iterable[int]]r   r  rq  Optional[bool]rE   r  r   zCallable[[T], JITFunction[T]]r4   )rU  zOptional[T]r   r  r  r  r   r  r   r  rq  r  rE   r  r   z4Union[JITFunction[T], Callable[[T], JITFunction[T]]])5
__future__r   r   rX   r&   r;   rt   rv  r  r  collectionsr   	functoolsr   typingr   r   r	   r
   r   r   r   r   r   r   runtime.driverr   typesr   rI   r   r?   r   NodeVisitorr   r   r   r   r   r   r   r   r  r   r   r  r   rF   r  r  r  r  r  r   r2   r1   <module>r     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	2 &A$%q!&N@/!$ N@l
 
 
 
 #*.15>B #
 
 (	

 /
 %<
 
 
 #
 

 4 #*.15>B #44 	4
 (4 /4 %<4 4 4 :4x 0? ?DA!r2   