
    ^j              	         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
Z
d dlm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 d dlmZ d dlmZ d	d
lm Z  ddl!m!Z! ddlm"Z" d	dl#m$Z$m%Z%m&Z&m'Z' ddl(m)Z) d dl*m+Z+m,Z,m-Z- dZ.dZ/ ed      Z0 G d dejb                        Z2d6dZ3 G d d      Z4d7dZ5 G d dee0         Z6d Z7d Z8d Z9 G d d       Z:i Z;e G d! d"             Z<d# Z=d$ Z> G d% d&e:e6e0         Z?ed8d'       Z@edddddddd(	 	 	 	 	 	 	 	 	 	 	 	 	 d9d)       Z@	 d:dddddddd(	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d;d*Z@ G d+ d,      ZA G d- d.      ZBd/ ZCd0 ZD G d1 d2e:      ZE G d3 d4e:      ZFd5 ZGy)<    )annotationsdivisionN)defaultdict)	dataclass)cached_property)	CallableGenericIterableOptionalTypeVaroverloadDictAnyTuple)BaseBackend)
ModuleType   )knobs   )driver)_async_compile)find_paths_ifget_iterable_pathtype_canonicalisation_dictis_namedtuple)get_cache_key)get_cache_invalidating_env_varsnative_specialize_implirztriton.languagez"triton.experimental.gluon.languageTc                  ~     e Zd ZdZd fdZed        Zd Zd Zd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| _	        t        t        ddh| _        i | _        d| _        y )Nutf-8>
   intlenmaxminlistfloatprintrangegetattr
isinstancecopymathF)super__init__namehashlibsha256encodehasherglobals	nonlocalssupported_python_builtinsGLUON_MODULETRITON_MODULEsupported_modulesused_global_valsvisiting_arg_default_value)selfr3   r8   r9   src	__class__s        ]/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/triton/runtime/jit.pyr2   zDependenciesFinder.__init__.   st    	nnSZZ%89 "*
& 	"
" TV*/'    c                6    | j                   j                         S N)r7   	hexdigestr@   s    rC   retzDependenciesFinder.retY   s    {{$$&&rD   c                    t        j                  |j                        ryt        |dd      }|j	                  t
              S )NT
__module__ )inspect	isbuiltinfuncr-   
startswithr<   )r@   noderO   modules       rC   _is_triton_builtinz%DependenciesFinder._is_triton_builtin]   s6    TYY'|R0  //rD   c                >   t        |t              sJ | 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 )
NGlobal variable z has value z when compiling z, but inner kernel z has conflicting value z7 from when it was first compiled.  This is not allowed.noinlineFr$   )r.   JITCallabler>   keysRuntimeErrorr3   __name__update	cache_keystrr-   r7   r6   )r@   rO   kvar_name_v1v2func_keys           rC   _update_hashzDependenciesFinder._update_hashc   s/   $,,, &&++-0E0E0J0J0LL 	AKHa))!,EB))!,EBRx"&xjB4?OPTPYPY{Zmnrn{n{m|  }T  UW  TX  XO  P 	 	$$T%:%:;>>Cj%8998??734rD   c                    ddl m} |t        |      t        u ry t	        |dd      r#|j
                  D ]  }| j                  |        y t	        |dd      ry t	        |dd      dk(  ry t        |t              r| j                  |       y t        |      r*t        |t              st        ||      st        d	|       | j                  ry |/t        j                  |      |f| j                  |t!        |      f<   y )
Nr   	constexpr__triton_aggregate__F__triton_builtin__rK   rL   ztriton.language.extra.libdevicez!Unsupported function referenced: )language.corerg   typer   r-   
hash_attrsrecord_referencer.   rW   rd   callablerY   r?   r/   deepcopyr>   id)r@   valvar_dictr3   rg   attrs         rC   rm   z#DependenciesFinder.record_referenceu   s    - ;$s)z13.6 ,%%d+,3,e4 3b)-NNc;'c"C=C!6z#y?Y!B3%HII **;?==;Mx:XD!!4H"67rD   c                >    t        |j                        t        j                  u r|j                  S |j                   j
                  v ry  fd} ||j                        \  }}|j                   j                  v r|S  j                  |||j                         |S )Nc                    j                   j                  | d       }||j                   fS j                  j                  | d       }||j                  fS y)NNN)r8   getr9   )r3   rq   r@   s     rC   name_lookupz2DependenciesFinder.visit_Name.<locals>.name_lookup   sZ    ,,""4.CDLL((..$$T40CDNN**rD   )rk   ctxastStorerp   local_namesr:   rm   )r@   rQ   rx   rq   rr   s   `    rC   
visit_NamezDependenciesFinder.visit_Name   s    >SYY&77N77d&&&	 $DGG,X77d444Jc8TWW5
rD   c                ^    |j                   D cg c]  }| j                  |       c}S c c}w rF   )eltsvisit)r@   rQ   elts      rC   visit_TuplezDependenciesFinder.visit_Tuple   s$     ,0995C

3555s   *c                f   | j                  |j                        }t        |t        j                        r6| j                  |j                        }t        |t        j                        r6t        |dd      }||| j                  v ry t        ||j                        }| j                  |       |S )NrZ   rL   )	r   valuer.   rz   	Attributer-   r=   rs   rm   )r@   rQ   lhslhs_namerI   s        rC   visit_Attributez"DependenciesFinder.visit_Attribute   s    jj$cmm,**SYY'C cmm,3
B/;(d&<&<<c499%c"
rD   c                    |j                   j                   D ch c]  }|j                   c}| _        | j                  |       y c c}w rF   )argsargr|   generic_visit)r@   rQ   r   s      rC   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?   r   )defaultsexprr@   s     rC   visit_defaultsz:DependenciesFinder.visit_arguments.<locals>.visit_defaults   sS    8::::26/$ )D'

4() 38/%/s   < < 	A)
	itertoolschainposonlyargsr   vararg
kwonlyargsr   kw_defaultskwargr   )r@   rQ   r   r   s   `   rC   visit_argumentsz"DependenciesFinder.visit_arguments   s    	8 ??4#3#3TYYQUQ\Q\bdfjfufuv 	CJJsO	 	t''(::!JJtzz"t}}%rD   c                    | j                  |      }t        |t              r| xj                  t	        |      z  c_        y | j                  j                  |       y rF   )r   r.   r)   r|   setadd)r@   rQ   targets      rC   visitAssnTargetz"DependenciesFinder.visitAssnTarget   sE     D!fd#F+  (rD   c                    t        |j                        dk7  rt        d      | j                  |j                  d          | j	                  |       y )Nr   z2Simultaneous multiple assignment is not supported.r   )r&   targets	TypeErrorr   r   r@   rQ   s     rC   visit_AssignzDependenciesFinder.visit_Assign   sG    t||!
 PQQT\\!_- 	4 rD   c                \    | j                  |j                         | j                  |       y rF   r   r   r   r   s     rC   visit_AnnAssignz"DependenciesFinder.visit_AnnAssign   $    T[[) 	4 rD   c                \    | j                  |j                         | j                  |       y rF   r   r   s     rC   	visit_ForzDependenciesFinder.visit_For  r   rD   )returnNonerv   )rZ   rK   __qualname____doc__r2   propertyrI   rS   rd   rm   r}   r   r   r   r   r   r   r   r   __classcell__rB   s   @rC   r"   r"   "   s`    	)0V ' '05$%N06
	!
&@)!!!rD   r"   c                    dd l mc m} t        | t              r| j                         } | j                  d      r7| j                  d      } t        |       } | j                  d      sJ d| dd  z   S | j                  d      rdt        | d d       z   S | j                  d      rdt        | dd        z   S | j                  d      rt        | j                  d            S t        | |j                        rdt        | j                         S t        | |j                        r| j                  } n(t        | t              r| j                  } nt	        |       } t!        j"                  | j%                  d	d
      |       S )Nr   zconst const**kr   ztl._trL   )triton.language.corelanguagecorer.   r]   striprP   removeprefix_normalize_tyendswithpointer_type
element_tydtyper3   rk   rZ   r   rw   replace)tyr   s     rC   r   r     s?   ''"cXXZ==")Br"B==%%%"QR&= ;;sr#2w///==r!"v...== !788	B))	*=/011	B

	#WW	B	[[W%))"**T2*>CCrD   c                      e Zd ZdZ	 	 ddZed        Zedd       Zed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 rF   )num_paramdo_not_specializedo_not_specialize_on_alignment)r@   r   paramr   r   s        rC   r2   zKernelParam.__init__.  s     !2.L+rD   c                .    | j                   j                  S rF   )r   r3   rH   s    rC   r3   zKernelParam.name5  s    {{rD   c                    | j                   j                  r1| j                   j                  t        j                  j                  k(  ryt        | j                   j                        S )NrL   )r   
annotationrM   	Parameteremptyr   rH   s    rC   r   zKernelParam.annotation9  sD    {{%%)?)?7CTCTCZCZ)ZT[[3344rD   c                    | j                   }|j                  d      r|dd  }n|j                  d      r|dd  }|t        t        j                               v r| j                   S y)Nr   r   r   r   rL   )r   rP   r   r   values)r@   as     rC   annotation_typezKernelParam.annotation_type?  s]    OO<<!"A\\#!"A.55788??"rD   c                    d| j                   v S Nrg   )r   rH   s    rC   is_constexprzKernelParam.is_constexprJ  s    doo--rD   c                r    | j                   ryd| j                  v xs | j                  j                  d      S )NFr   r   )r   r   rP   rH   s    rC   is_constzKernelParam.is_constN  s1    $//)MT__-G-G-MMrD   c                .    | j                   j                  S rF   )r   defaultrH   s    rC   r   zKernelParam.defaultT  s    {{"""rD   c                d    | j                   j                  t        j                  j                  k7  S rF   )r   r   rM   r   r   rH   s    rC   has_defaultzKernelParam.has_defaultX  s#    {{""g&7&7&=&===rD   N)r   r%   r   zinspect.Parameterr   boolr   r   r   r]   )rZ   rK   r   r   r2   r   r3   r   r   r   r   r   r   r    rD   rC   r   r   +  s    LM15M     5 5
   . . N N
 # # > >rD   r   c                6    d}d}t        t        | |||      d   S )NFTr   )r   r   )r   
specializer   aligns       rC   mangle_typer   ]  s%    HE!+sHj%PQRSSrD   c                  ,    e Zd ZU ded<   d Zd ZddZy)KernelInterfacer    runc               \     | j                   t        t        j                  |      |dd|S )NTgridwarmup)r   map
MockTensor
wrap_dtype)r@   r   r   kwargss       rC   r   zKernelInterface.warmupf  s*    txxZ5J5JD1QT$\U[\\rD   c                   t        d      )Nzrun not implemented)NotImplementedError)r@   r   r   r   r   s        rC   r   zKernelInterface.runi  s    !"788rD   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 )NFr   )r   )r   r   r   r@   s     rC   <lambda>z-KernelInterface.__getitem__.<locals>.<lambda>r  s    xtxx$T%'YRX'Y rD   r   )r@   r   s   ``rC   __getitem__zKernelInterface.__getitem__l  s     ZYrD   N)r   r    )rZ   rK   r   __annotations__r   r   r   r   rD   rC   r   r   c  s    	
F]9ZrD   r   c           
        |j                         D ci c]  \  }}||j                  j                  dk(  rt        |      nh|j                  j                  dk(  rd|j                  inA|j                  j                  dk(  r'd|j
                   d|j                  j                   in| }}}dd l}| ||j                         D 	cg c]  }	t        |	       c}	t        |j                               |j                         D 	cg c]  }	t        |	       c}	t        |j                               |j                  |j                  d	}
|j                  |
      }|S c c}}w c c}	w c c}	w )Nr   rg   JITFunctionjit_function:r   )	r3   	signatureconstant_keysconstant_vals
attrs_keys
attrs_valsoptionskeyr   )itemsrB   rZ   r]   r   rR   fnr   jsonrX   r)   r   __dict__dumps)r3   r   	constantsattrsr   r  r   r   r  xobjserialized_objs               rC   serialize_specialization_datar  v  sI   
 $//+	 C 	5??33w>SZ??##{2 FQRWR]R]D^??##}4 :HELL>YZ[`[c[c[p[pZqIr8s:?	@I  9QZQ_Q_Qa?bAQ?bY %**,0OQa0O_cdidpdpdr_s##C6??C
 ZZ_N @c0Os   BEE$E)c                    t        | j                        t        |      k(  sJ g }t        | j                  j                         |      D ]  \  }}|j                  r|j                  d| d       )|j                  rdnd}|j                  rdnd}|j                  rdnd}d| d| d| d| d	}	|j                  rt        |j                  t              r"|j                  dk(  s|j                  dd	 d
v rd}|r#|j                  d|j                   d|	 d       |j                  d|j                   d       |j                  |	        	 d }
ddj                  t        t        |
| j                  j                                     dgz          ddj                  | j                  j                         D cg c]
  }d| d|  c}       ddj                  |       d}| j                  j                         D ci c];  \  }}|j                   t"        j$                  j&                  urd| |j                   = }}}t(        }||d<   ||d<   t*        |d<   t-        ||       |d   S 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("constexpr", )TrueFalsezspecialize_impl(backend, , u1Nr   )fpbfFz("z",) + z[1:]z", None)c                x    | d   j                   t        j                  j                  u r| d   S | d    d| d    S )Nr   r   z	=default_)r   rM   r   r   )r	  s    rC   r   z0create_function_from_signature.<locals>.<lambda>  sA    AaDLLG,=,=,C,CCAaD AaD6QZ[\]^[_Z`Ia rD   z
def dynamic_func(z	**optionsz):
    params = {'z': z}
    specialization = [,z-]
    return params, specialization, options
default_specialize_implbackendrW   dynamic_func)r&   
parametersziprX   r   appendr   r   r   r   r.   r]   joinr)   r   r  r   rM   r   r   r   rW   exec)sigkparamsr  specializationr3   kpr   r   r   rI   r   	func_bodyr   func_namespacer  s                  rC   create_function_from_signaturer'    s    s~~#g,...N++-w7 0b??!!N4&":;!#v'H$&$8$8fJ!@@GfE-dV2hZr*RPUwVWXC!!b00#6))T1R5G5G5K|5[%*
"))Br/A/A.B&T*RS #))Br/A/A.B(*KL%%/'0, bC))DS#..*>*>*@!ABk]RST U		3>>;N;N;PQ4QtfCv.QRS Txx/0 1I >>//1D%== 1 1 7 77 4&5==(N  -O(7N$% 'N9$/N=! 	N# .))+ Rs    I5A I:c                8    | j                    d| j                   S )N.)rK   r   r  s    rC   get_full_namer+    s    mm_Aboo.//rD   c                  t    e Zd Zd Zd Zedd       Zd Zd Zed        Z	ddZ
d Zd	 Zd
 Z eee      Zy)rW   c                   || _         t        j                  |      | _        	 t        j                  |      \  | _        | _        t        |      | _	        t        j                         | _        t        j                  dj                  | j                              }|t!        j"                  d|t         j$                        j'                         d  }|| _        d | _        i | _        |j.                  | _        |j0                  | _        |j2                  | _        |j4                  | _        |j6                  | _        y # t        $ r}t        d      |d }~ww xY w)Nz1@jit functions should be defined in a Python filerL   z^def\s+\w+\s*\()r  rM   r   getsourcelinesraw_srcstarting_line_numberOSError
ValueErrorr+  _fn_name	threadingRLock
_hash_locktextwrapdedentr  research	MULTILINEstart_srchashr>   r   rZ   r   __globals__rK   )r@   r  erA   s       rC   r2   zJITCallable.__init__  s    **2.	Y6=6L6LR6P3DL$3 &b)#//+ oobggdll34")).R\\BHHJKL		 TV zzOO>>--7  	YPQWXX	Ys   "D= =	EEEc                   | j                   }|j                  | j                  S t        |j                  j
                  |j                        D ci c]  \  }}||j                   }}}| j                  |z  S c c}}w rF   )r  __closure__r?  r  __code__co_freevarscell_contents)r@   r  r3   cellr9   s        rC   get_capture_scopezJITCallable.get_capture_scope  sr    WW>>!###@CBKKD[D[]_]k]k@lm*$T4---m	m)++ ns   A;c                   | j                   5  | j                  | j                  cd d d        S d| j                   | _        t        j                  | j
                        j                  }t        | j                  | j                  || j                        }|j                  | j                                |j                  t        | j                        z   | _        t        t!        |j"                  j%                                     | _        ddlm} | xj                  t        | j"                  j%                         D cg c]  \  \  }}\  }}t+        ||      r||f c}}}      z  c_        t-        j.                  | j                  j1                  d            j3                         | _        d d d        | j                  S c c}}}w # 1 sw Y   | j                  S xY w)Nz
recursion:)r3   r8   r9   rA   r   rf   r$   )r6  r>  r3  rM   getclosurevarsr  r9   r"   r?  rA   r   parserI   r]   r0  dictsortedr>   r  r   rg   r.   r4   r5   r6   rG   )r@   r9   dependencies_finderrg   r3   r`   rq   s          rC   r\   zJITCallable.cache_key  s    __ 	Nyy$yy	N 	N
 %T]]O4DI..tww7AAI"4$--QUQaQamv9=#C%%djjl3+//#d6O6O2PPDI$(0C0T0T0Z0Z0\)]$^D!6II9=9N9N9T9T9V= ="5)4Xc1!+C!;  $Sk = > >I  tyy'7'7'@AKKMDI#	N$ yy	=	N$ yys$   GDG G
%AG
GG%c                ,    t        | j                        S rF   )r>  r\   rH   s    rC   __hash__zJITCallable.__hash__
  s    DNN##rD   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   rJ  r=  r.   Moduler&   bodyFunctionDef)r@   trees     rC   rJ  zJITCallable.parse  s_    yy#$

+++499~"""$))A,888rD   c                    ddl m}  ||       S )Nr   )constexpr_type)r   rV  )r@   rV  s     rC   rk   zJITCallable.type  s    7d##rD   c                     y rF   r   )r@   handless     rC   _flatten_irzJITCallable._flatten_ir  s    rD   c                     d| _         || _        y)a"  
        The only method allowed to modify src.
        Bypasses the __setattr__ restriction by calling super().__setattr__ directly.

        Note that it is the callers responsibility to make sure any triton functions that call this function have the `.hash` value reset to None.
        N)r>  r=  )r@   new_srcs     rC   _unsafe_update_srczJITCallable._unsafe_update_src  s     		rD   c                    t        d      )NzqCannot set attribute 'src' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)AttributeErrorrH   s    rC   _set_srczJITCallable._set_src)  s     ( ) 	)rD   c                    | j                   S rF   )r=  rH   s    rC   _get_srczJITCallable._get_src.  s    yyrD   )fgetfsetNr   )rX  zlist[ir.value]r   r   )rZ   rK   r   r2   rG  r   r\   rO  rJ  rk   rY  r\  r_  ra  rA   r   rD   rC   rW   rW     sa     (D,  ,$ $ $)
 x
0CrD   rW   c                  ,    e Zd ZU ded<   ded<   ded<   y)JitFunctionInfor   rR   r]   r3   r   r   N)rZ   rK   r   r   r   rD   rC   re  re  7  s    
IrD   re  c                    t        |      t        |      f}| j                  |d       }||S fdt         |            t        |      z   }|| |<   |S )Nc                F   t        | t              r| D cg c]
  } |       c}S t        |       r$| D cg c]
  } |       }} | j                  | S t        | t              rt	        fd| D              S t        | t
              r| j                  S | S c c}w c c}w )Nc              3  .   K   | ]  } |        y wrF   r   ).0r   replace_callabless     rC   	<genexpr>z?compute_cache_key.<locals>.replace_callables.<locals>.<genexpr>L  s     ?C*3/?s   )r.   r)   r   rB   tuplerW   r\   )r
  r   resultsrj  s      rC   rj  z,compute_cache_key.<locals>.replace_callablesE  s    c4 69:s%c*::39<=#(-=G= 3=='**U#?3???[)== 
 ;=s
   BB)rl  r]   rw   )kernel_key_cacher#  r   r  r\   rj  s        @rC   compute_cache_keyro  >  se     #g,
/C $$S$/I
 %n56WEI%SrD   c                ~    t        | t              s| S t        |       D ]  \  }}t        |      | |<    t	        |       S rF   )r.   r)   	enumerateconvert_to_tuple_if_listrl  )iteminested_values      rC   rr  rr  V  sD    dD! %T? 9<*<8Q9 ;rD   c                  l     e Zd Zd Z	 	 ddZd Zd Zd Zd Zd Z		 	 d fd	Z
d	 Zd
 Zd Zd Z xZS )r   c                     y)NFr   rH   s    rC   is_gluonzJITFunction.is_gluond  s    rD   c
                   |s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|j                   d	| d
}t        | j                         }t        ||||d   |||      }||||j                  |j                  |j                  |j                  |j                  |j                  |||	d} |||t        ||
|       d|i||	d      S c c}}w )Nr  r   z: z[num_warps=z, num_ctas=z, num_stages=z, enable_fp_fusion=z, launch_cooperative_grid=](r  r   )r   devicer  	num_warpsnum_ctas
num_stagesenable_fp_fusionlaunch_cooperative_gridextern_libsconfigsspecialization_data	is_warmupr  F)r  reprr  compileis_manual_warmupalready_compiled)r  r   rK   r  r  paramsr3   r|  r}  r~  r  r  r+  r  r  re  )r@   hookr  r   r   r{  r  r   r  r  r3   rR   r   r   	arg_reprsr  	full_namer  r   s                      rC   
_call_hookzJITFunction._call_hookg  s    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  |[  |[  {\  \^  _h  ^i  ij  k!$''*	;IyR[]def]gipru<BD #" **((!,, ' 8 8'.'F'F"..#6"
 vtT2C*6*&"
 	
- `s   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)rn   pre_run_hooksr  )r@   r  s     rC   add_pre_run_hookzJITFunction.add_pre_run_hook  s$    
 ~~!!$'rD   c                    ddl m}m}m}m} t
        j                  j                         } ||      }|| _        || _        || _        t        | j                  | j                  |      }i i |||fS )z1
        Precompute as much as possible.
        r   )CompiledKernelr  	ASTSourcemake_backend)compilerr  r  r  r  r   activeget_current_targetr'  r   r  )r@   r  r  r  r  r   r  binders           rC   create_binderzJITFunction.create_binder  sd     	PO113v&,"/WU2vw..rD   c           
        |j                  |      }| j                  D cg c]  }|j                   }}|D cg c]  }|d   	 }}t        ||      D 	
ci c]  \  }	}
|	|

 }}	}
d|vsJ d       d|vsJ d       d|vsJ d       |D ]#  }	|	|j                  vs|	|vst        d|	z         t        |d	       }|D ci c]&  }|t        t        |j                               |      ( }}|D cg c]  }|d   d
k(  rdn|d    }}t        |d       }|D 	ci c]  }	|	|j                  t        ||	              }}	||||fS c c}w c c}w c c}
}	w c c}w c c}w c c}	w )Nr   device_typez=device_type option is deprecated; current target will be usedr{  z8device option is deprecated; current device will be usedstreamz8stream option is deprecated; current stream will be usedz2Keyword argument %s was specified but unrecognisedc                    |dk(  S r   r   )r`   rq   s     rC   r   z(JITFunction._pack_args.<locals>.<lambda>  s    3+;M rD   rg   rL   r   c                "    t        |t              S rF   )r.   r]   )r`   r	  s     rC   r   z(JITFunction._pack_args.<locals>.<lambda>  s    Z35G rD   )parse_optionsr  r3   r  r  KeyErrorr   r   r)   r   
parse_attr)r@   r  r   
bound_argsr#  r   r	  sigkeyssigvalsr^   vr   
constexprspathattrvalsr  s                   rC   
_pack_argszJITFunction._pack_args  s   ''/#';;/a166//!/0A1Q400(+GW(=>fq!QT>	>F*k,kk*v%a'aa%v%a'aa% 	YA(((Qg-=SVWWXX	Y #7,MN
[efSWd-d:3D3D3F.GNNf
fAOPA!A$+-B1Q47PPh(GHPUV1G&&'81'EFFVV	:u44% 00> gPVs"   EE
E +E1E#Ec               H   |j                  d| j                        xs t        j                  j                  |d<   t        j                  j
                  |d<   t        j                  j                         }t        j                  j                  |      }| j                  D ]
  } ||i |  | j                  |   \  }}	}
}} ||i |\  }}}t        j                  j                  6t        j                  j                         \  }}|j                  d| d       t        |	||      }|j                  |d       }|4| j                  |||||      \  }}}}| j!                  |||||||      }|y t#               }| j$                  j'                         D ]6  \  \  }}\  }}|j                  ||      x}|k7  s$t)        d| d| d|        |s|J t+        |      r ||      }t-        |      }|d   }|d	kD  r|d	   nd	}|d
kD  r|d
   nd	}  |j.                  ||g|j1                          }! |j2                  ||| ||j4                  |j6                  |!t        j                  j8                  t        j                  j:                  g	|j1                           |S )Ndebuginstrumentation_modez("custom_pipeline", r  rU   z1 has changed since we compiled this kernel, from z to r   r   r   )rw   r  r   runtimecompilationr  r   r  get_current_deviceget_current_streamr  device_cachesadd_stages_inspection_hookr  ro  r  _do_compileobjectr>   r  rY   rn   r&   launch_metadatar   r   functionpacked_metadatalaunch_enter_hooklaunch_exit_hook)"r@   r   r   r   r   r{  r  r  kernel_cachern  r   r  r  r  r#  r   inspect_stages_keyinspect_stages_hashr  kernelr   r  r  not_presentr3   r`   rq   globals_dictnewVal	grid_sizegrid_0grid_1grid_2r  s"                                     rC   r   zJITFunction.run  s    **Wdjj9PU]]=P=Pw).):):)O)O%& 11311&9 && 	"D$!&!	" CGBTBTU[B\?& /5d.Ef.E+
NG ==33?6;mm6^6^6`3 3!!$89L8MQ"OP 0.'J!!#t, >48OOGVU_aoDK5M1GY
E %%c9fj'SXZ`aF~ 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4f44T6XJDUDUDWXOFJJvvvvvH^H^`o}}668V8VnYcYjYjYlnrD   c                T    | j                   | j                  S | j                  |      S rF   )_reprr3  )r@   r`   s     rC   r  zJITFunction.repr  s"     $

 2t}}E

1ErD   c	           	     P   |r|ng }|r|ng }t         |   |       |j                  | _        || _        || _        || _        || _        || _        | t        | j                   d| j                  j                   <   g | _        t        | j                  j                  j!                               D ]T  \  }	}
|	|v xs |
j"                  |v }|	|v xs |
j"                  |v }| j                  j%                  t'        |	|
||             V t)        | j*                        | _        d | _        || _        || _        | j                  D cg c]  }|j"                   c}| _        | j                  D cg c]  }|j6                  s|j8                   c}| _        g | _        y c c}w c c}w )Nr   )r1   r2   rK   rR   versionr   r   r  r  _triton_jit_function_registryr  r   r  rq  r   r  r   r3   r  r   r   r  r  r  r  rV   	arg_namesr   r   r  r  )r@   r  r  r   r   r  rV   r  r  rt  r   dnsdns_oaprB   s                 rC   r2   zJITFunction.__init__   s|   1B-Ki)Goq&mm!2.L+
.QU%Qtww7K7K6L&MN!$..";";"B"B"DE 	CHAu((KEJJ:K,KC88hEJJJh<hFKK{1eS&AB	C )););< 
  +/++6Q!&&6*.++HQ155H  	 7Hs   F.F# F#c           	        dd l }dd lm t        j                  j                         }|j                  |      }|d   | j                  k7  rt        d|d    d| j                         t        t        |d         }|d   }| j                  |   \  }}}}	}|d   }
|j                  |
k7  rt        d|
 d|       fd}t        ||      D ci c]  \  }}| ||       }}}t        t        |d	         }|d
   }t        t        ||            }|d   j                         D ci c]  \  }}|t!        |       }}}|d   j                         D ci c]#  \  }}|t#        |t$              rt        |      n|% }}}|d   }|	j'                  |      }| j)                  ||||||d      S c c}}w c c}}w c c}}w )Nr   r3   zSpecialization data is for z but trying to preload for r   r   r   c                   j                   j                  |       rj                  |       S t        | t              rAd| v rj	                  | d         S d| v r%| d   }|t
        v r	t
        |   S t        d| d      | S )Nrg   r   zUnable to resolve JITFunction z for preload)r   is_dtyper.   rK  rg   r  rY   )r   jf_keytls     rC   _decode_constantz-JITFunction.preload.<locals>._decode_constant7  s    xx  'xx&%&%'<<k(:;;!U*">2F!>><VDD&)Gx|'\]]LrD   r   r   r   r   r  T)r   )r  triton.languager   r   r  r  loadsr3  rY   r   rl  r  r  r  rK  r  rr  r.   r)   r  r  )r@   r  r  r{  deserialized_objr   r   r`   r   r  deserialized_targetr  r  r   r  r   r   r  r   r   r  s                       @rC   preloadzJITFunction.preload&  s   $113::&9:F#t}}4-.>v.F-GGbcgcpcpbqrt tE#3O#DE(9#'#5#5f#= 1fgq.x8 ??11!<=P<QQlmsltuvv	 FIXeEfgzsEc+E22g
g 0 >?
%l3
SZ01 M]]hLiLoLoLqrjc5S2599r	r /y9??A
U E4!8ueC
 
 u%''0   
 	
 h s
s   F2*F8(F>c                R     j                      \  }}	} j                  t        j                  j                  g	      ry  j                         t        j                  j                         }
|
It               t        |	      } fd} f
d}|
j                  |||      }|<   |S  j                  j                        }|<    j                  t        j                  j                  g	       |S )Nc                 B    j                  j                         S )N)r   r   	_env_vars)r  r  )env_varsr   r@   rA   r   s   rC   async_compilez.JITFunction._do_compile.<locals>.async_compilei  s!    ||C@P@P\d|eerD   c                t   
 | <   j                  t        j                  j                  	g
	       y rF   )r  r   r  jit_post_compile_hook)r  r  r  r{  r  r  r   r@   r   r   r   s    rC   finalize_compilez1JITFunction._do_compile.<locals>.finalize_compilel  s;    $*S! C CS)U[]ceo '%&:rD   )r   r   )r  r  r   r  jit_cache_hookr  r   active_moderw   r   r   submitr  r  r  )r@   r  r   r{  r  r   r  r   r`   r  
async_moder\   r  r  r  r  r  rA   r   s   ````````       @@@@rC   r  zJITFunction._do_compile[  s0   .2.@.@.H+a!??5==77iQWYcelotnu!#nnT9j%@#//335
!68H%c7GXFIf f: :
  &&y-AQRF &L 	 \\#fg>N>N\OF &LOOEMM??iQWY_akmt"GV-rD   c                    t        d      )Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rY   r@   r   r   s      rC   __call__zJITFunction.__call__z  s    WXXrD   c                P    d| j                    d| j                  j                   dS )NzJITFunction(r   r  )rR   r  r   rH   s    rC   __repr__zJITFunction.__repr__}  s&    dkk]!DGG,@,@+ACCrD   )r   zbool | None)NNNNNNN)rZ   rK   r   rx  r  r  r  r  r   r  r2   r  r  r  r  r   r   s   @rC   r   r   b  sZ    .
 
.
`(/507rF mq;?$ L3
j>YDrD   r   c                     y rF   r   r*  s    rC   jitr    s    rD   r  r  r  r   r   r  rV   c                     y rF   r   r  s          rC   r  r    s     rD   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                  j                  rddlm}  ||       S t        |       S )Nr   )InterpretedFunction)r  r   r   r  rV   r  r  )rn   r   r  	interpretinterpreterr  r   )	r  r  r  r   r   r  rV   r  r  s	     rC   	decoratorzjit.<locals>.decorator  sj    ||==""8&r7N_Fdlq08tUdf f "3/M! /	 	rD   r  r    r   zJITFunction[T]r   )	r  r  r  r  r   r   r  rV   r  s	    ``````` rC   r  r    s&    : & 
~} rD   c                  N    e Zd ZdZed        Zd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)rB   rZ   rK   r   )r   s    rC   r   zMockTensor.wrap_dtype  s.    ==!!W,71Jc?"
rD   Nc                *    |dg}|| _         || _        y )Nr   )r   shape)r@   r   r  s      rC   r2   zMockTensor.__init__  s    =CE

rD   c                    dg}| j                   dd  D ]  }|j                  |d   |z          t        t        |            S )Nr   r   )r  r  rl  reversed)r@   stridessizes      rC   stridezMockTensor.stride  sG    #JJqrN 	/DNN72;-.	/Xg&''rD   c                      yNr   r   r   rD   rC   data_ptrzMockTensor.data_ptr      rD   c                      yr   r   r   rD   rC   	ptr_rangezMockTensor.ptr_range  r  rD   rF   )
rZ   rK   r   r   staticmethodr   r2   r  r  r  r   rD   rC   r   r     sM    
  
(    rD   r   c                  J    e Zd Zd Zd Zd ZddZd Zd Zd Z	d Z
d	 Zd
 Zy)TensorWrapperc                    || _         || _        |j                  | _        |j                  | _        | j                  j                  | _        y rF   )r   basedatar{  r  )r@   r	  r   s      rC   r2   zTensorWrapper.__init__  s5    
	II	kkYY__
rD   c                6    | j                   j                         S rF   )r	  r  rH   s    rC   r  zTensorWrapper.data_ptr  s    yy!!##rD   c                4     | j                   j                  | S rF   )r	  r  )r@   r   s     rC   r  zTensorWrapper.stride  s    tyy&&rD   c                <    d| j                    d| j                   dS )NzTensorWrapper[rz  r  )r   r	  rH   s    rC   __str__zTensorWrapper.__str__  s    

|2dii[::rD   c                6    | j                   j                         S rF   )r	  element_sizerH   s    rC   r  zTensorWrapper.element_size  s    yy%%''rD   c                ^    t        | j                  j                         | j                        S rF   )r  r	  cpur   rH   s    rC   r  zTensorWrapper.cpu  s    TYY]]_djj99rD   c                N    | j                   j                  |j                          y rF   )r	  copy_)r@   others     rC   r  zTensorWrapper.copy_  s    		

#rD   c                ^    t        | j                  j                         | j                        S rF   )r  r	  cloner   rH   s    rC   r  zTensorWrapper.clone  s    TYY__.

;;rD   c                `    t        | j                  j                  |      | j                        S rF   )r  r	  tor   )r@   r{  s     rC   r  zTensorWrapper.to  s     TYY\\&14::>>rD   c                `    t        | j                  j                  |      | j                        S rF   )r  r	  	new_emptyr   )r@   sizess     rC   r  zTensorWrapper.new_empty  s"    TYY007DDrD   Nr   )rZ   rK   r   r2   r  r  r  r  r  r  r  r  r  r   rD   rC   r  r    s5    %$';(:$<?ErD   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   hasattrr   rk   )tensorr   s     rC   reinterpretr     sk    &-(FKK%%%;; !e44		$VU++/V~Q?@@rD   c                L   | }t        |t              s|j                  }t        |t              s|j                  j                  j                  }|j
                  }t        |j                        D ].  \  }}|j                         j                  d      s&||z  } ||fS  ||fS )Nzdef )
r.   rW   r  rC  co_filenamer0  rq  r/  r   rP   )r  base_fn	file_name
begin_lineidxlines         rC   get_jit_fn_file_liner(  *  s    G+.** +.

##//I--J w/ 	T::<""6*#Jj  	 j  rD   c                  (    e Zd Zd Zed        Zd Zy)BoundConstexprFunctionc                     || _         || _        y rF   )__self____func__)r@   instancer  s      rC   r2   zBoundConstexprFunction.__init__>  s     rD   c                .    | j                   j                  S rF   )r-  r\   rH   s    rC   r\   z BoundConstexprFunction.cache_keyB  s    }}&&&rD   c                B     | j                   | j                  g|i |S rF   )r-  r,  r  s      rC   r  zBoundConstexprFunction.__call__F  s!    t}}T]]<T<V<<rD   N)rZ   rK   r   r2   r   r\   r  r   rD   rC   r*  r*  <  s      ' '=rD   r*  c                  0     e Zd Z fdZd ZdddZ xZS )ConstexprFunctionc                $    t         |   |       y rF   )r1   r2   )r@   r  rB   s     rC   r2   zConstexprFunction.__init__L  s    rD   c                "    |t        ||       S | S rF   )r*  )r@   r
  objclasss      rC   __get__zConstexprFunction.__get__O  s    ?)#t44rD   N)	_semanticc                  ddl m}m} |D cg c]
  } ||       }}|j                         D ci c]  \  }}| ||       }}} | j                  |i |}	||	S t
        j                  j                  r|	S  ||	      S c c}w c c}}w )Nr   )_unwrap_if_constexprrg   )r   r9  rg   r  r  r   r  r  )
r@   r7  r   r   r9  rg   r	  r^   r  ress
             rC   r  zConstexprFunction.__call__U  s    H156A$Q'66;A<<>J!Q!)!,,JJ dggt&v&J ==""J~ 7Js
   BB)rZ   rK   r   r2   r6  r  r   r   s   @rC   r2  r2  J  s     )-  rD   r2  c                    t        |       S )z
    Wraps an arbitrary Python function so that it can be called at
    compile-time on constexpr arguments in a Triton function and
    returns a constexpr result.
    )r2  r*  s    rC   constexpr_functionr<  h  s     R  rD   r   )Fr  )r  Optional[Callable]r  r=  r   Optional[Iterable[int | str]]r   r>  r  Optional[bool]rV   r?  r   zCallable[[T], JITFunction[T]]rF   )r  zOptional[T]r  r=  r  r=  r   r>  r   r>  r  r?  rV   r?  r   zKernelInterface[T])H
__future__r   r   rz   r/   r4   rM   r   r4  r9  r7  collectionsr   dataclassesr   	functoolsr   typingr   r	   r
   r   r   r   r   r   r   triton.backendsr   typesr   rL   r   r   r   _utilsr   r   r   r   cacher   triton._C.libtritonr   r   r   r<   r;   r    NodeVisitorr"   r   r   r   r   r  r'  r+  rW   r  re  ro  rr  r   r  r   r  r   r(  r*  r2  r<  r   rD   rC   <module>rK     s;   , 
      	  # ! % ] ] ] '     ` `   [ [!3CLg! g!^D4/> />dTZgaj Z&$9*x0i1 i1X !#    0	\D+q1 \DH	 
 
 
 #*.7;DH #
 
 (	

 5
 %B
 
 
 #
 

 4 #*.7;DH #44 	4
 (4 54 %B4 4 4 4x B"E "EJA!$=[ = <!rD   