
    wi              	         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j!                  Z"dWdZ# G d d          Z$d Z%i Z&dXd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/d0d1d2d3d4d5d6d7d8d9d:d;d<Z+ e,e+-                                          D ]Z.e.e+e.<    G d= d>e(e                    Z/edYdA            Z0eddddddddBdZdM            Z0	 d[ddddddddBd\dPZ0 G dQ dR          Z1 G dS dT          Z2dU Z3dV Z4dS )]    )annotationsdivisionN)defaultdict)cached_property)
CallableGenericIterableOptionalTypeVarUnionoverloadDictAnyTuple   )driver)
ModuleTypez.runtime.jitTc                  |     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).
    returnNonec                    t                                                       || _        t          j        |                    d                    | _        || _        h d| _        i | _	        d| _
        d S )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       b/root/.openclaw/workspace/chatterbox_venv_py311/lib/python3.11/site-packages/triton/runtime/jit.pyr&   zDependenciesFinder.__init__$   st    	nSZZ%8%899 *
 *
 *
&. TV*/'''    c                4    | j                                         S N)r+   	hexdigestr0   s    r3   retzDependenciesFinder.retH   s    {$$&&&r4   c                    t          j        |j                  rdS t          |dd          }|                    t
                    S )NT
__module__ )inspect	isbuiltinfuncr#   
startswithTRITON_MODULE)r0   noder?   modules       r3   _is_triton_builtinz%DependenciesFinder._is_triton_builtinL   sA    TY'' 	4|R00  ///r4   c                2   t          |t                    r | j                                        |j                                        z  D ]V}|\  }}| j        |         \  }}|j        |         \  }}||k    r)t	          d| d| d| j         d|j         d| d          W| j                            |j                   |j        }|t          t          |dd                    z  }| j                            |                    d	                     d S d S )
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*   )r0   r?   kvar_name_v1v2func_keys           r3   _update_hashzDependenciesFinder._update_hashR   se   dK(( 	9 *//11D4I4N4N4P4PP  !-a0A-a0A88& T8  T  T  T  TTXT]  T  Trvr  T  T  Y[  T  T  T    !(()>???~HGD*e<<===HKxw7788888	9 	9r4   c                   t          |j                  t          j        u r|j        S |j        | j        v rd S | j                            |j        d           }|{| j        stt          |          t          ur^t          |t                    sIt          |dd          s8|j        | j        vr*|| j        f| j        |j        t	          | j                  f<   |                     |           |S )N__triton_builtin__F)typectxastStoreidlocal_namesr,   getr/   r   r$   rH   r#   r-   r.   rU   )r0   rB   vals      r3   
visit_NamezDependenciesFinder.visit_Named   s    >>SY&&7N7d&&&4ltw--
 O 7	  IIZ// #344 0 >ESJ^`e=f=f 0 G4#AAABEt|ATD!47Bt|,<,<"=>#
r4   c                *      fd|j         D             S )Nc                :    g | ]}                     |          S  )visit).0eltr0   s     r3   
<listcomp>z2DependenciesFinder.visit_Tuple.<locals>.<listcomp>   s#    555C

3555r4   )eltsr0   rB   s   ` r3   visit_TuplezDependenciesFinder.visit_Tuple   s!     6555495555r4   c                f   |                      |j                  }t          |t          j                  r4|                      |j                  }t          |t          j                  4|t          |dd          t          k    rd S t          ||j                  }|                     |           |S )NrK   r<   )	rd   valuer$   rZ   	Attributer#   rA   attrrU   )r0   rB   lhsr9   s       r3   visit_Attributez"DependenciesFinder.visit_Attribute   s    jj$$cm,, 	(**SY''C cm,, 	(;73
B77=HH4c49%%#
r4   c                f    d |j         j         D             | _        |                     |           d S )Nc                    h | ]	}|j         
S rc   arg)re   rt   s     r3   	<setcomp>z7DependenciesFinder.visit_FunctionDef.<locals>.<setcomp>   s    >>>CG>>>r4   )argsr]   generic_visitri   s     r3   visit_FunctionDefz$DependenciesFinder.visit_FunctionDef   s6    >>ty~>>>4     r4   c                .     fd}t          j        |j        |j        |j        r|j        gng |j                  D ]}                     |            ||j                   |j                             |j                    ||j	                   d S )Nc                    	 j         rJ d_         | D ]}|                    |           	 d_         d S # d_         w xY w)NTF)r/   rd   )defaultsexprr0   s     r3   visit_defaultsz:DependenciesFinder.visit_arguments.<locals>.visit_defaults   sp    8::::26/$ ) )D'

4((() 38///%/7777s	   ,9 	A)
	itertoolschainposonlyargsrv   vararg
kwonlyargsrd   kw_defaultskwargr{   )r0   rB   r}   rt   s   `   r3   visit_argumentsz"DependenciesFinder.visit_arguments   s    	8 	8 	8 	8 	8 ?4#3TYQUQ\@dbdfjfuvv 	 	CJJsOOOOt'(((:!JJtz"""t}%%%%%r4   c                    |                      |          }t          |t                    r| xj        t	          |          z  c_        d S | j                            |           d S r6   )rd   r$   r   r]   setadd)r0   rB   targets      r3   visitAssnTargetz"DependenciesFinder.visitAssnTarget   sd     D!!fd## 	)F+  (((((r4   c                    t          |j                  dk    rt          d          |                     |j        d                    |                     |           d S )N   z2Simultaneous multiple assignment is not supported.r   )r   targets	TypeErrorr   rw   ri   s     r3   visit_AssignzDependenciesFinder.visit_Assign   s^    t|!!
 PQQQT\!_--- 	4     r4   c                d    |                      |j                   |                     |           d S r6   r   r   rw   ri   s     r3   visit_AnnAssignz"DependenciesFinder.visit_AnnAssign   4    T[))) 	4     r4   c                d    |                      |j                   |                     |           d S r6   r   ri   s     r3   	visit_ForzDependenciesFinder.visit_For   r   r4   )r   r   )rK   r;   __qualname____doc__r&   propertyr9   rD   rU   r`   rj   rp   rx   r   r   r   r   r   __classcell__r2   s   @r3   r   r      s	       	 	"0 "0 "0 "0 "0 "0H ' ' X'0 0 09 9 9$  <6 6 6
  ! ! !
& & &@) ) )! ! !! ! !! ! ! ! ! ! !r4   r   r   rN   c                    t          | t                    r| j        S t          | t                    r| S t	          |           S r6   )r$   rX   rK   rN   repr)tys    r3   _normalize_tyr      s>    "d {	B		 	88Or4   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dS )KernelParamzBRepresents a parameter (name plus metadata) to a @jit'ed function.numr   paraminspect.Parameterdo_not_specializebooldo_not_specialize_on_alignmentc                >    || _         || _        || _        || _        d S r6   )r   _paramr   r   )r0   r   r   r   r   s        r3   r&   zKernelParam.__init__   s&    !2.L+++r4   c                    | j         j        S r6   )r   r'   r8   s    r3   r'   zKernelParam.name   s    {r4   c                    | j         j        r| j         j        t          j        j        k    rdS t          | j         j                  S )Nr<   )r   
annotationr=   	Parameteremptyr   r8   s    r3   r   zKernelParam.annotation   s<    {% 	)?7CTCZ)Z)Z2T[3444r4   c                    | j         }dD ]@\  }}||                    |          t          |          z   d          }|r||v r| | c S A|dk    rdS dS )N))uintu)r   ir   u1r<   )r   findr   )r0   r   ty1ty2widths        r3   annotation_typezKernelParam.annotation_type   s    _
5 	' 	'HCzs33c#hh>??@E '
**u&&&4rr4   c                    d| j         v S )N	constexpr)r   r8   s    r3   is_constexprzKernelParam.is_constexpr  s    do--r4   c                $    d| j         v o| j         S )Nconst)r   r   r8   s    r3   is_constzKernelParam.is_const	  s    $/)C$2C.CCr4   c                    | j         j        S r6   )r   defaultr8   s    r3   r   zKernelParam.default  s    {""r4   c                @    | j         j        t          j        j        k    S r6   )r   r   r=   r   r   r8   s    r3   has_defaultzKernelParam.has_default  s    {"g&7&===r4   N)r   r   r   r   r   r   r   r   )rK   r;   r   r   r&   r   r'   r   r   r   r   r   r   r   rc   r4   r3   r   r      s        LLM M M M     _  5 5 _5
   _ . . _. D D _D # # X# > > X> > >r4   r   c                    |r-t          | d          r|                                 dz  dk    rdS t          | t                    r|r| dz  dk    rdS | dk    rdS dS )Ndata_ptr   r   Dr   1N)hasattrr   r$   r   )valigns     r3   compute_spec_keyr     su     J'' QZZ\\B->!-C-Cs	As		  	a"fkk3!VV33r4   Fc                   | dS t          | t                    rdS t          | t                    rd| k    r| dk    rdS d| k    r| dk    rdS d	S t          | t                    rd
S t	          | d          rdS | j        |f}t                              |d           }|P|d         rdndt          t          |d                   
                    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_dictrN   split)rt   r   dskress       r3   mangle_typer   &  s
   
{v	C		 t	C		 s??si//5c\\cY..55	C		 v	(	)	) 	{ y(#mmC&&;q6*44s.HSQRVIZIZ[^I_I_`bIc.ddC IcN
r4   c                  "    e Zd ZU ded<   ddZdS )KernelInterfacer   runr   c                      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   )rv   kwargsr   r0   s     r3   <lambda>z-KernelInterface.__getitem__.<locals>.<lambda>J  s     xtx$T%'Y'YRX'Y'Y r4   rc   )r0   r   s   ``r3   __getitem__zKernelInterface.__getitem__D  s     ZYYYYYr4   N)r   r   )rK   r;   r   __annotations__r   rc   r4   r3   r   r   A  s9         
FFFZ Z Z Z Z Zr4   r   c                    d |                                 D             }dd l}| |||                                |j        |d}|                    |          }|S )Nc                X    i | ]'\  }}||j         j        d k    rt          |          n|(S r   )r2   rK   rN   re   keyrl   s      r3   
<dictcomp>z1serialize_specialization_data.<locals>.<dictcomp>O  s:    wwwWaWZ\aEO$<$G$Gc%jjjUwwwr4   r   )r'   	signature	constantsattrsoptionsr   )itemsjsonto_dict__dict__dumps)	r'   r   r   r   r   r   r   objserialized_objs	            r3   serialize_specialization_datar  N  sg    wwenetetevevwwwIKKK99u}} C ZZ__Nr4   c                @   t          | j                  t          |          k    sJ g }g }g }g }g }g }t          | j                                        |          D ]C\  \  }	}
}|
j        t
          j        j        u r1|                    |	           |                    d|	 d|	            n5|                    |	 d|	            |                    d|	 d|	            |j	        r|                    |	           |                    |	           |j
        s8|j        s|                    d|	z             n|                    d|	z             |j        r|                    d|j        z             |                    d|	d|j        rd	nd
d           Ed                    d ||z   D                       }d                    d |D                       }d                    d |D                       }|                    d           d                    |          }d                    |          }d|d|d|d|d|d}d | j                                        D             }t          |d<   |j        |d<   t#          ||           |d         S )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<   c                    g | ]}|d z   S r  rc   re   xs     r3   rg   z2create_function_from_signature.<locals>.<listcomp>  s    MMMaTMMMr4   c                    g | ]}|d z   S r	  rc   r
  s     r3   rg   z2create_function_from_signature.<locals>.<listcomp>  s    ???1a$h???r4   c                    g | ]}|d z   S r	  rc   r
  s     r3   rg   z2create_function_from_signature.<locals>.<listcomp>  s    !G!G!Gq!d(!G!G!Gr4   z**excess_kwargszdef dynamic_func(z):
    return {z}, (z), (z), excess_kwargsc                Z    i | ](\  }}|j         t          j        j        ud | |j         )S )default_)r   r=   r   r   )re   r'   r   s      r3   r   z2create_function_from_signature.<locals>.<dictcomp>  sE       D%= 1 777 	45=777r4   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kprM   args_strdict_str	func_bodyfunc_namespaces                    r3   create_function_from_signaturer%  Y  s(    s~#g,,.... ILNOO 4 4 6 6@@ k k$R:*000T""" 3D 3 3T 3 3444455t55666 3D 3 3T 3 3444? 	k!!$''''%%d+++' W8 W#**+MPT+TUUUU#**+NQU+UVVV! k&&v0B'BCCCC&&&PRP[FhffahFhFh'ijjjjMM?_+LMMMNNIWW?????@@N!G!G4F!G!G!GHH&''' yy##Hyy&&HH(((III~~~7I7I7IKI >//11  N %0N=!)0)AN%& 	N### .))r4   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 )rH   Nc                >   t          | d          r| j        S t          | t                    rdS t          | t                    rd| k    r| dk    rdS d| k    r| dk    rdS d	S t          | t
                    rd
S | d S 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   rX   rs   s    r3   _key_ofzJITFunction._key_of  s    3   	G9T"" 	G4S!! 	G33)#3#3u##"2"2uuU## 	G6[4ES		EEEEFFFr4   Fc                    | dS t          | t                    r| S t          |                               d          d         }t          |         }|rdnd}||z   S )N*i8r   r   r   r   )r$   rN   r   r   )r   r   	dtype_str	const_strs       r3   _type_ofzJITFunction._type_of  sd     ;5S!! 	JHHNN3''+	.y9	$-DD#	9$$r4   c                J    t          t          | j        |                    }|S r6   )dictr  
constexprs)r0   constexpr_keyr   s      r3   _make_constantszJITFunction._make_constants  s!    T_m<<==	r4   c	                
   |rt           j        nt           j        }	|	dS | j        j        }
| j        j        }d                    d t          | j        |d                   D                       }|
 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 )NFr  c                ,    g | ]\  }}|j          d | S )z: r'   )re   r   r   s      r3   rg   z*JITFunction._call_hook.<locals>.<listcomp>  s,    ___%*4444___r4   r   z[num_warps=z, num_ctas=z, num_stages=z, enable_fp_fusion=](r  c                      e Zd Zd ZdS )/JITFunction._call_hook.<locals>.JitFunctionInfoc                0    || _         || _        || _        d S r6   )rC   r'   jit_function)r0   rC   r'   r\  s       r3   r&   z8JITFunction._call_hook.<locals>.JitFunctionInfo.__init__  s    $ 	$0!r4   N)rK   r;   r   r&   rc   r4   r3   JitFunctionInforZ    s#            r4   r]  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)rH   
cache_hookcompiled_hookrg  rK   r;   r  r  paramsr_  r`  ra  rb  r  rc  )r0   r   r   r^  r   r   rd  rf  beforehookr'   rC   	arg_reprsr   r]  re  r   s                    r3   
_call_hookzJITFunction._call_hook  s    *0N{%%[5N<5w#II__c$+WZ[\W]F^F^___``	  p  p7#4  p  pAQ  p  p`g`r  p  p  HO  H`  p  p  dm  p  p  p	 	 	 	 	 	 	 	 <D)YX_`aXbdkmpqq #" *(!, ' 8".#6"
 
 tvtT22C*6*&"
 
 
 	
r4   c                \    t          |          sJ | j                            |           dS )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  )r0   ro  s     r3   add_pre_run_hookzJITFunction.add_pre_run_hook  s2    
 ~~~!!$'''''r4   c                h   ddl m}m}m}m} || _        || _        || _        || _        t          | j        | j        |          | _        d t          | j                  D             | _
        d t          | j                  D             | _        d t          | j                  D             | _        dS )z1
        Precompute as much as possible.
        r   )CompiledKernelrh  	ASTSourcemake_backendc                &    g | ]\  }}|j         |S rc   r   re   r   ps      r3   rg   z-JITFunction.create_binder.<locals>.<listcomp>-  s#    ![![![AAN![!![![![r4   c                &    g | ]\  }}|j         |S rc   r{  r|  s      r3   rg   z-JITFunction.create_binder.<locals>.<listcomp>.  s%    %c%c%cFQTUTb%ca%c%c%cr4   c                4    g | ]\  }}|j         |j        |S rc   )r   r   r|  s      r3   rg   z-JITFunction.create_binder.<locals>.<listcomp>/  s?     $
 $
 $
1a1;N$
YZYg$
$
 $
 $
r4   N)compilerrw  rh  rx  ry  r%  r   rm  binder	enumerateconstexpr_indicesnon_constexpr_indicesspecialised_indices)r0   r  rw  rh  rx  ry  s         r3   create_binderzJITFunction.create_binder#  s     	POOOOOOOOOOO,"(4T^T[RYZZ![![)DK2H2H![![![%c%ci6L6L%c%c%c"$
 $
%dk22$
 $
 $
   r4   c               2   ( |                     dd          p#t          j                             dd          dk    |d<   ddlm} t
          j                                        }t
          j                            |          }t
          j        	                                } ||          }	 j
        D ]
}
 |
|i |  j                             |	             j        |i |\  }}}}}d                    |          t          ||f          z   } j        |                              |d           }||	                    |          }d	|vs
J d
            d|vs
J d            d|vs
J d            |D ]}||j        vrt%          d|z            t'          |                                          } fd j        D             }|d t-          |                   }d t/          ||          D             }|	                     j        |          f}|d                                         ((fdt/          | j                  D             }|                                D ]'\  }}t9          |          rt;          d| d          (                     |||||||d          rd S                       |||d                   }                      |||j                  }| j        |         |<                        |||||||d           tC                      } j"                                        D ]?\  \  }}\  } }!|!                     ||          x}"| k    rtG          d| d|  d|"           @|s|J t9          |          r ||          }t-          |          }#|d         }$|#dk    r|d         nd}%|#dk    r|d         nd}& |j$        ||g|R  }' |j%        |$|%|&||j&        |j'        |' j(        j)         j(        j*        g	|R   |S )NdebugFTRITON_DEBUG0r   r   )ry  r<   device_typez=device_type option is deprecated; current target will be usedr^  z8device option is deprecated; current device will be usedstreamz8stream option is deprecated; current stream will be usedz2Keyword argument %s was specified but unrecognisedc                4    g | ]}j         |         j        S rc   )rm  r'   )re   r   r0   s     r3   rg   z#JITFunction.run.<locals>.<listcomp>\  s"    OOOqt{1~*OOOr4   c                *    i | ]\  }}||d k    rdn|S )r   rL  rc   )re   rO   r   s      r3   r   z#JITFunction.run.<locals>.<dictcomp>^  s*    ```AqF{{UU```r4   r   c                J    i | ]\  }}|j         s|j        v s||j        | S r6   )r   r   r'   )re   r   r}  constant_paramss      r3   r   z#JITFunction.run.<locals>.<dictcomp>b  sF       Q> '(e&>&>19 CL99r4   zCallable constexpr at index z is not supportedT)rn  )r   r   rF   z1 has changed since we compiled this kernel, from z to r   )+r^   osenvironr  ry  r   activeget_current_deviceget_current_streamget_current_targetrt  r  r  r  rN   cacheparse_optionsr   KeyErrortuplevaluesr  r   r  get_attrs_descriptorrm  get_constantsr   rs  r   rq  rx  rh  objectr.   rJ   launch_metadatar   functionpacked_metadatarw  launch_enter_hooklaunch_exit_hook))r0   r   r   rv   r   ry  r^  r  r   r  ro  
bound_argssig_and_specr  r  excess_kwargsr   kernelr   rO   
bound_valssigkeyssigvalsr   rd  r   r   rt   r1   not_presentr'   rQ   r_   globals_dictnewVal	grid_sizegrid_0grid_1grid_2r  r  s)   `                                       @r3   r   zJITFunction.run3  s    **We44b
~WZ8[8[_b8bw 	,+++++113311&991133,v&& & 	" 	"DD$!&!!!!;w'''VaVZVacgVrkqVrVrS
L.2Dm ggl##c>=*I&J&JJF#''T22>++F33G !...0o...6)))+e)))6)))+e)))" ] ]G,,,"#WZ[#[\\\ - z002233J POOOD4NOOOG"=CLL=1G``#gW^J_J_```I33DKLLOG%aj6688O   !*dk::  I
 $//++ Y Y3C== Y#$W1$W$W$WXXXY sIvy'7TZcghh t..y)WQZHHC\\( "  F
 '-DJvs#OOCFIwQW`eOfff hh.2.C.I.I.K.K 	q 	q*IT1*\&**4===#EE"otoo^aoogmooq q q F  	y###~~ ( tJ''D		I!WF )AT!WW1F )AT!WW1F 5f4T6WDVWWWOFJvvvvvH^`o*<d>Q>byewy y y yr4   c	           	     $   |r|ng }|r|ng }| _         j        | _        || _        t	          j                  | _        || _        || _        t	          j                  d         | _	        fd| _
        || _        d | _        g | _        t          | j        j                                                  D ]I\  }	}
|	|v p|
j        |v }|	|v p|
j        |v }| j                            t'          |	|
||                     Jt)          j        t	          j                            | _        | j        t1          j        d| j        t0          j                                                  d          | _        t9          t:                    | _        d | _        i | _         d | _!        || _"        d | j        D             | _#        d | j        D             | _$        g | _%        j&        | _&        j'        | _'        j(        | _(        j        | _        d S )Nr   c                ,    j         n
 |           S r6   )rK   )rQ   rg  r   s    r3   r   z&JITFunction.__init__.<locals>.<lambda>  s    T\bkkttAww r4   z^def\s+\w+\s*\(c                    g | ]	}|j         
S rc   rW  re   r}  s     r3   rg   z(JITFunction.__init__.<locals>.<listcomp>  s    666Q!&666r4   c                *    g | ]}|j         	|j        S rc   )r   r   r  s     r3   rg   z(JITFunction.__init__.<locals>.<listcomp>  s!    HHHQH15HHHr4   ))rg  r;   rC   versionr=   r   r   r   getsourcelinesstarting_line_numberr   r  r  rm  r  r  r  r'   r  r   textwrapdedent	getsourcer1   research	MULTILINEstartr   rQ  r  hashr.   r  rG   	arg_namesrR  rt  r   rK   __globals__)r0   rg  r  r   r   r  rG   r   r  r   r   dnsdns_oas    `     `     r3   r&   zJITFunction.__init__  s   1BJ--Ki)q)G)Goq&m *2..!2.L+$+$:2$>$>q$A!FFFFF	.!$.";"B"B"D"DEE 	C 	CHAu((KEJ:K,KC88hEJJh<hFK{1eS&AABBBB ?7#4R#8#8998BI&8$(BLQQWWYYZZ[ &&
	 TV   76$+666HH$+HHH   z>-r4   c                `   | j         t          | j        | j        | j                  }|                    |                                            |j        t          | j	                  z   | _         t          t          |j                                                            | _        | j         S )N)r'   r,   r1   )r  r   rK   r  r1   rd   parser9   rN   r  rQ  sortedr.   r   )r0   dependencies_finders     r3   rM   zJITFunction.cache_key  s     9"4$-QUQagkgo"p"p"p%%djjll333+/#d6O2P2PPDI$(0C0T0Z0Z0\0\)])]$^$^D!yr4   c               R     | j         t          t          j        |          |dd|S )NTr   )r   map
MockTensor
wrap_dtype)r0   r   rv   r   s       r3   r   zJITFunction.warmup  s.    txZ5JD1Q1QT$\\U[\\\r4   c           	        ddl m}m} ddlm} dd l}dd lm t          j	        
                                }|                    |          }|d         | j        j        k    r%t          d|d          d| j        j                   fd|d	                                         D             }t!          |d
                                                   }	 || |	||                    |d                             }
d |d                                         D             }|d         } ||
d |          }|| j        |         |<   |S )Nr   )rh  rx  r   )AttrsDescriptorr'   zSpecialization data is for z but trying to preload for c                z    i | ]7\  }}|j                             |          r                     |          n|8S rc   )r   is_dtype)re   r   rl   tls      r3   r   z'JITFunction.preload.<locals>.<dictcomp>  sR     
 
 
U BH$5$5e$<$<G%%
 
 
r4   r   r   r   c                b    i | ],\  }}|t          |t                    rt          |          n|-S rc   )r$   r   r  r   s      r3   r   z'JITFunction.preload.<locals>.<dictcomp>  sG     
 
 
U E4!8!8Cue
 
 
r4   r   r   )r  rh  rx  triton.backends.compilerr  r   triton.languagelanguager   r  r  loadsrg  rK   rJ   r   rQ  	from_dictr  )r0   re  rh  rx  r  r   r^  deserialized_objr   r   r1   r   r   r  r  s                 @r3   preloadzJITFunction.preload  s   11111111<<<<<<$$$$$$1133::&9::F#tw'777u.>v.Fuucgcjcsuuw w w
 
 
 
.{;AACC
 
 
	 )+6<<>>??	iiO4M4MN^_fNg4h4hii
 
.y9??AA
 
 
 u%dG,,"(
63r4   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   )rZ   r  r1   r$   Moduler   bodyFunctionDef)r0   trees     r3   r  zJITFunction.parse  sg    y""$
+++++49~~""""$)A,88888r4   c                     t          d          )Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rJ   )r0   rv   r   s      r3   __call__zJITFunction.__call__  s    WXXXr4   c                x    t          t          |                               ||           |dk    r	d | _        d S d S )Nr1   )r%   rH   __setattr__r  )r0   r'   rl   r2   s      r3   r  zJITFunction.__setattr__   s@    k4  ,,T5999 5==DIII =r4   c                2    d| j          d| j        j         dS )NzJITFunction(:r  )rC   rg  rK   r8   s    r3   __repr__zJITFunction.__repr__  s"    ?dk??DG,<????r4   F)NNNNNNN)rK   r;   r   rk  rl  staticmethodrJ  rO  rT  rq  ru  r  r   r&   r   rM   r   r  r  r  r  r  r   r   s   @r3   rH   rH     sb       J MG G \G& 
% 
% 
% \
%  3
 3
 3
j( ( (
 
 
 \ \ \| mq;?:( :( :( :(x   X] ] ]  8  Y Y Y    @ @ @ @ @ @ @r4   rH   rg  JITFunction[T]c                    d S r6   rc   )rg  s    r3   jitr    s    Cr4   r  r   r  r   r   r  rG   r   Optional[Callable]r  r   Optional[Iterable[int]]r   r  Optional[bool]rG   Callable[[T], JITFunction[T]]c                    d S r6   rc   r  s          r3   r  r    s	     Cr4   Optional[T]4Union[JITFunction[T], Callable[[T], JITFunction[T]]]c               F    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
    rg  r   r   r  c           
         t          |           sJ t          j        dd          dk    rddlm}  ||           S t          |           S )NTRITON_INTERPRETr  r   r   )InterpretedFunction)r  r   r   r  rG   r   r  )rs  r  getenvinterpreterr  rH   )	rg  r  r  r   r   r  rG   r   r  s	     r3   	decoratorzjit.<locals>.decorator@  s    |||9'--44888888&&r7N_Fdlq08tUdf f f f "3/M! /	 	 	 	r4   Nrg  r   r   r  rc   )	rg  r  r   r  r   r   r  rG   r  s	    ``````` r3   r  r  #  sb    :           & 
~y}} r4   c                  Z    e Zd ZdZed             Zd Zed             Zed             ZdS )r  zr
    Can be used in place of real tensors when calling:
        kernel.warmup(MockTensor(torch.float32), ...)
    c                Z    | j         j        dk    r| j        dk    rt          |           S | S )Nr   torch)r2   rK   r;   r  rs   s    r3   r  zMockTensor.wrap_dtypee  s/    =!W,,71J1Jc??"
r4   c                    || _         d S r6   r   )r0   r   s     r3   r&   zMockTensor.__init__k  s    


r4   c                     dS Nr   rc   rc   r4   r3   r   zMockTensor.data_ptrn      qr4   c                     dS r  rc   rc   r4   r3   	ptr_rangezMockTensor.ptr_ranger  r  r4   N)	rK   r;   r   r   r  r  r&   r   r  rc   r4   r3   r  r  _  s~         
   \
     \   \  r4   r  c                  F    e Zd Zd Zd Zd ZddZd Zd Zd	 Z	d
 Z
d ZdS )TensorWrapperc                t    || _         || _        |j        | _        |j        | _        | j        j        | _        d S r6   )r   basedatar^  shape)r0   r  r   s      r3   r&   zTensorWrapper.__init__y  s1    
	I	kY_


r4   c                4    | j                                         S r6   )r  r   r8   s    r3   r   zTensorWrapper.data_ptr  s    y!!###r4   c                6    | j                             |          S r6   )r  stride)r0   r   s     r3   r  zTensorWrapper.stride  s    y"""r4   r   rN   c                (    d| j          d| j         dS )NzTensorWrapper[rX  r  )r   r  r8   s    r3   __str__zTensorWrapper.__str__  s    :
::di::::r4   c                4    | j                                         S r6   )r  element_sizer8   s    r3   r  zTensorWrapper.element_size  s    y%%'''r4   c                Z    t          | j                                        | j                  S r6   )r  r  cpur   r8   s    r3   r  zTensorWrapper.cpu  s    TY]]__dj999r4   c                D    | j                             |j                    d S r6   )r  copy_)r0   others     r3   r  zTensorWrapper.copy_  s    	
#####r4   c                Z    t          | j                                        | j                  S r6   )r  r  cloner   r8   s    r3   r  zTensorWrapper.clone  s     TY__..
;;;r4   c                \    t          | j                            |          | j                  S r6   )r  r  tor   )r0   r^  s     r3   r  zTensorWrapper.to  s"    TY\\&114:>>>r4   Nr   rN   )rK   r;   r   r&   r   r  r  r  r  r  r  r  rc   r4   r3   r  r  w  s        % % %$ $ $# # #; ; ; ;( ( (: : :$ $ $< < <? ? ? ? ?r4   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   rX   )tensorr   s     r3   reinterpretr    s    &-(( AFK%%%; !e444		$	$ AVU+++?V???@@@r4   c                P   | }t          |t                    s|j        }t          |t                    |j        j        j        }t          j        |j                  \  }}t          |          D ]3\  }}|                                	                    d          r||z  } n4||fS )Nzdef )
r$   rH   rg  __code__co_filenamer=   r  r  stripr@   )rg  base_fn	file_namelines
begin_lineidxlines          r3   get_jit_fn_file_liner)    s    G+.. * +.. 
#/I.wz::E: u%%  	T::<<""6** 	#JE	 j  r4   r  r  r  )r   r  r  r  r   r  r   r  r  r  rG   r  r   r  r6   )rg  r  r   r  r  r  r   r  r   r  r  r  rG   r  r   r  )5
__future__r   r   rZ   r(   r=   r~   r  r  r  collectionsr   	functoolsr   typingr   r   r	   r
   r   r   r   r   r   r   runtime.driverr   typesr   rK   r   rA   r   NodeVisitorr   r   r   r   r   r   r   r  r%  r   r   r  r   rH   r  r  r  r  r)  rc   r4   r3   <module>r1     s   , , , , , , , , 



       				 				  # # # # # # % % % % % % d d d d d d d d d d d d d d d d d d d d d d d d # # # # # #      .33~..../GCLL~! ~! ~! ~! ~! ~! ~! ~!L   -> -> -> -> -> -> -> ->`
 
 
 	   6	Z 	Z 	Z 	Z 	Zgaj 	Z 	Z 	Z  A* A* A*H
D)  :	
 Y ) y 7 : z v  v v D  U!" U#$ -   2 
(//11	2	2 & &A$%q!!N@ N@ N@ N@ N@/!$ N@ N@ N@l
 
   
 
 #*.15>B #
 
 
 
 
 

 4 #*.15>B #4 4 4 4 4 4x       0? ? ? ? ? ? ? ?DA A A! ! ! ! !r4   