
    Iri|              	         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, dZ-dZ. ed          Z/ G d dej0                  Z1dGdZ2 G d d          Z3dHdZ4 G d dee/                   Z5d Z6d  Z7d! Z8 G d" d#          Z9e G d$ d%                      Z:d& Z;d' Z< G d( d)e9e5e/                   Z=edId,            Z>edddddddd-dJd8            Z>	 dKdddddddd-dLd;Z> G d< d=          Z? G d> d?          Z@d@ ZAdA ZB G dB dCe9          ZC G dD dEe9          ZDdF ZEdS )M    )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_implz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).
    returnNonec                "   t                                                       || _        t          j        |                    d                    | _        || _        || _        h d| _	        t          t          ddh| _        i | _        d| _        d S )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)selfr4   r9   r:   src	__class__s        j/var/www/html/bestrading.cuttalo.com/models/btc_v9/venv/lib/python3.11/site-packages/triton/runtime/jit.pyr3   zDependenciesFinder.__init__.   s    	nSZZ%8%899 "*
 *
 *
& 	"
" TV*/'''    c                4    | j                                         S N)r8   	hexdigestrA   s    rD   retzDependenciesFinder.retY   s    {$$&&&rE   c                    t          j        |j                  rdS t          |dd          }|                    t
                    S )NT
__module__ )inspect	isbuiltinfuncr.   
startswithr=   )rA   noderP   modules       rD   _is_triton_builtinz%DependenciesFinder._is_triton_builtin]   sA    TY'' 	4|R00  ///rE   c                0   t          |t                    sJ | 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 )
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RuntimeErrorr4   __name__update	cache_keystrr.   r8   r7   )rA   rP   kvar_name_v1v2func_keys           rD   _update_hashzDependenciesFinder._update_hashc   s_   $,,,,, &++--0E0J0J0L0LL 	 	AKHa)!,EB)!,EBRxx" Px  P  PB  P  PPTPY  P  Pnrn{  P  P  UW  P  P  P    	$$T%:;;;>Cj%889998??73344444rE   Nc                R   ddl m} |t          |          t          u rd S t	          |dd          r!|j        D ]}|                     |           d S t	          |dd          rd S t	          |dd          dk    rd S t          |t                    r| 	                    |           d S t          |          r7t          |t                    s"t          ||          st          d	|           | j        rd S |-t          j        |          |f| j        |t!          |          f<   d S )
Nr   	constexpr__triton_aggregate__F__triton_builtin__rL   rM   ztriton.language.extra.libdevicez!Unsupported function referenced: )language.corerh   typer   r.   
hash_attrsrecord_referencer/   rX   re   callablerZ   r@   r0   deepcopyr?   id)rA   valvar_dictr4   rh   attrs         rD   rn   z#DependenciesFinder.record_referenceu   sd   ------ ;$s))z11F3.66 	 , ,%%d++++F3,e44 	F 3b))-NNNFc;'' 	c"""FC== 	JC!6!6 	Jz#y?Y?Y 	JH3HHIII * 	F;?=;M;Mx:XD!4H"67rE   c                     t          |j                  t          j        u r|j        S |j         j        v rd S  fd} ||j                  \  }}|j         j        v r|S                      |||j                   |S )Nc                    j                             | d           }|	|j         fS j                            | d           }|	|j        fS dS )NNN)r9   getr:   )r4   rr   rA   s     rD   name_lookupz2DependenciesFinder.visit_Name.<locals>.name_lookup   sZ    ,""4..CDL((.$$T400CDN**:rE   )rl   ctxastStorerq   local_namesr;   rn   )rA   rR   ry   rr   rs   s   `    rD   
visit_NamezDependenciesFinder.visit_Name   s    >>SY&&7N7d&&&4	 	 	 	 	 $DG,,X7d444Jc8TW555
rE   c                *      fd|j         D             S )Nc                :    g | ]}                     |          S  )visit).0eltrA   s     rD   
<listcomp>z2DependenciesFinder.visit_Tuple.<locals>.<listcomp>   s#    555C

3555rE   )eltsrA   rR   s   ` rD   visit_TuplezDependenciesFinder.visit_Tuple   s!     6555495555rE   c                f   |                      |j                  }t          |t          j                  r4|                      |j                  }t          |t          j                  4t          |dd          }|	|| j        v rd S t          ||j                  }|                     |           |S )Nr[   rM   )	r   valuer/   r{   	Attributer.   r>   rt   rn   )rA   rR   lhslhs_namerJ   s        rD   visit_Attributez"DependenciesFinder.visit_Attribute   s    jj$$cm,, 	(**SY''C cm,, 	(3
B//;(d&<<<4c49%%c"""
rE   c                f    d |j         j         D             | _        |                     |           d S )Nc                    h | ]	}|j         
S r   arg)r   r   s     rD   	<setcomp>z7DependenciesFinder.visit_FunctionDef.<locals>.<setcomp>   s    >>>CG>>>rE   )argsr}   generic_visitr   s     rD   visit_FunctionDefz$DependenciesFinder.visit_FunctionDef   s6    >>ty~>>>4     rE   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@   r   )defaultsexprrA   s     rD   visit_defaultsz:DependenciesFinder.visit_arguments.<locals>.visit_defaults   sp    8::::26/$ ) )D'

4((() 38///%/7777s	   ,9 	A)
	itertoolschainposonlyargsr   vararg
kwonlyargsr   kw_defaultskwargr   )rA   rR   r   r   s   `   rD   visit_argumentsz"DependenciesFinder.visit_arguments   s    	8 	8 	8 	8 	8 ?4#3TYQUQ\@dbdfjfuvv 	 	CJJsOOOOt'(((:!JJtz"""t}%%%%%rE   c                    |                      |          }t          |t                    r| xj        t	          |          z  c_        d S | j                            |           d S rG   )r   r/   r*   r}   setadd)rA   rR   targets      rD   visitAssnTargetz"DependenciesFinder.visitAssnTarget   sd     D!!fd## 	)F+  (((((rE   c                    t          |j                  dk    rt          d          |                     |j        d                    |                     |           d S )Nr   z2Simultaneous multiple assignment is not supported.r   )r'   targets	TypeErrorr   r   r   s     rD   visit_AssignzDependenciesFinder.visit_Assign   s^    t|!!
 PQQQT\!_--- 	4     rE   c                d    |                      |j                   |                     |           d S rG   r   r   r   r   s     rD   visit_AnnAssignz"DependenciesFinder.visit_AnnAssign   4    T[))) 	4     rE   c                d    |                      |j                   |                     |           d S rG   r   r   s     rD   	visit_ForzDependenciesFinder.visit_For  r   rE   )r"   r#   rw   )r[   rL   __qualname____doc__r3   propertyrJ   rT   re   rn   r~   r   r   r   r   r   r   r   r   __classcell__rC   s   @rD   r!   r!   "   s       	 	)0 )0 )0 )0 )0 )0V ' ' X'0 0 05 5 5$% % % %N  06 6 6
	 	 	! ! !
& & &@) ) )! ! !! ! !! ! ! ! ! ! !rE   r!   r"   r^   c                   dd l mc m} t          | t                    r|                                 } |                     d          rH|                     d          } t          |           } |                     d          sJ d| dd          z   S | 	                    d          rdt          | d d                   z   S |                     d          rdt          | dd                    z   S |                     d          r"t          |                     d                    S nut          | |j
                  rdt          | j                   S t          | |j                  r| j        } n,t          | t                    r| j        } nt	          |           } t!          j        |                     d	d
          |           S )Nr   zconst const**kr   ztl._trM   )triton.language.corelanguagecorer/   r^   striprQ   removeprefix_normalize_tyendswithpointer_type
element_tydtyper4   rl   r[   r   rx   replace)tyr   s     rD   r   r     s   '''''''''"c XXZZ=="" 	!))Br""B==%%%%%"QRR&= ;;s 	0r#2#w////== 	/r!""v....== 	9 !7!7888	9	B)	*	* 1=//111	B
	#	# W	B		 [WW%)"**T2*>*>CCCrE   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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 rG   )r   _paramr   r   )rA   r   r   r   r   s        rD   r3   zKernelParam.__init__.  s&    !2.L+++rE   c                    | j         j        S rG   )r   r4   rI   s    rD   r4   zKernelParam.name5  s    {rE   r"   r^   c                    | j         j        r| j         j        t          j        j        k    rdS t          | j         j                  S )NrM   )r   
annotationrN   	Parameteremptyr   rI   s    rD   r   zKernelParam.annotation9  s<    {% 	)?7CTCZ)Z)Z2T[3444rE   c                    | j         }|                    d          r|dd          }n|                    d          r
|dd          }|t          t          j                              v r| j         S dS )Nr   r   r   r   rM   )r   rQ   r   r   values)rA   as     rD   annotation_typezKernelParam.annotation_type?  sw    O<< 	!""AA\\# 	!""A.5778888?"rrE   c                    d| j         v S Nrh   )r   rI   s    rD   is_constexprzKernelParam.is_constexprJ  s    do--rE   c                Z    | j         rdS d| j        v p| j                            d          S )NFr   r   )r   r   rQ   rI   s    rD   is_constzKernelParam.is_constN  s4     	5$/)MT_-G-G-M-MMrE   c                    | j         j        S rG   )r   defaultrI   s    rD   r   zKernelParam.defaultT  s    {""rE   c                @    | j         j        t          j        j        k    S rG   )r   r   rN   r   r   rI   s    rD   has_defaultzKernelParam.has_defaultX  s    {"g&7&===rE   N)r   r&   r   r   r   r   r   r   r"   r^   )r[   rL   r   r   r3   r   r4   r   r   r   r   r   r   r   r   rE   rD   r   r   +  s        LLM M M M     _  5 5 5 _5
    _ . . _. N N _N
 # # X# > > X> > >rE   r   Fc                F    d}d}t          t          | |||          d         S )NFTr   )r   r   )r   
specializer   aligns       rD   mangle_typer   ]  s(    HE!+sHj%PPQRSSrE   c                  .    e Zd ZU ded<   d Zd ZddZdS )	KernelInterfacer   runc               R     | j         t          t          j        |          |dd|S )NTgridwarmup)r   map
MockTensor
wrap_dtype)rA   r   r   kwargss       rD   r   zKernelInterface.warmupf  s.    txZ5JD1Q1QT$\\U[\\\rE   c                    t          d          )Nzrun not implemented)NotImplementedError)rA   r   r   r   r   s        rD   r   zKernelInterface.runi  s    !"7888rE   r"   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   rA   s     rD   <lambda>z-KernelInterface.__getitem__.<locals>.<lambda>r  s     xtx$T%'Y'YRX'Y'Y rE   r   )rA   r   s   ``rD   __getitem__zKernelInterface.__getitem__l  s     ZYYYYYrE   N)r"   r   )r[   rL   r   __annotations__r   r   r   r   rE   rD   r   r   c  sZ         
FFF] ] ]9 9 9Z Z Z Z Z ZrE   r   c           	        d |                                 D             }dd l}| |d |                                D             t          |                                          d |                                D             t          |                                          |j        |d}|                    |          }|S )Nc                    i | ]@\  }}||j         j        d k    rt          |          n|j         j        dk    r	d|j        in|AS )r   rh   )rC   r[   r^   r   r   keyr   s      rD   
<dictcomp>z1serialize_specialization_data.<locals>.<dictcomp>w  sf        C 	5?3w>>SZZZ&+o&>+&M&Mek""SX  rE   r   c                ,    g | ]}t          |          S r   r*   r   xs     rD   r   z1serialize_specialization_data.<locals>.<listcomp>  s    ?b?b?bAQ?b?b?brE   c                ,    g | ]}t          |          S r   r  r  s     rD   r   z1serialize_specialization_data.<locals>.<listcomp>  s    0O0O0OQa0O0O0OrE   )r4   	signatureconstant_keysconstant_vals
attrs_keys
attrs_valsoptionsr   )itemsjsonrY   r*   r   __dict__dumps)	r4   r  	constantsattrsr  r   r  objserialized_objs	            rD   serialize_specialization_datar  v  s      $//++  I KKK9?b?bQZQ_Q_QaQa?b?b?bY  0O0O%**,,0O0O0O_cdidpdpdrdr_s_s#C C
 ZZ__NrE   c           
     j   t          | j                  t          |          k    sJ g }t          | j                                        |          D ]\  }}|j        r|                    d| d           &|j        rdnd}|j        rdnd}|j        rdnd}d| d| d| d| d	}	|j	        r{t          |j	        t                    r|j	        dk    s|j	        dd	         d
v rd}|r"|                    d|j	         d|	 d           |                    d|j	         d           |                    |	            d }
dd                    t          t          |
| j                                                            dgz              dd                    d | j                                        D                        dd                    |           d}d | j                                        D             }t           }||d<   ||d<   t"          |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("constexpr", )TrueFalsezspecialize_impl(backend, , u1Nr   )fpbfFz("z",) + z[1:]z", None)c                t    | d         j         t          j        j        u r| d         n| d          d| d          S )Nr   r   z	=default_r   rN   r   r   )r  s    rD   r   z0create_function_from_signature.<locals>.<lambda>  s=    AaDLG,=,CCCAaDDAaDIaIa[\]^[_IaIa rE   z
def dynamic_func(z	**optionsz):
    params = {c                    g | ]
}d | d| S )'z': r   )r   r4   s     rD   r   z2create_function_from_signature.<locals>.<listcomp>  s)    QQQ4.t....QQQrE   z}
    specialization = [,z-]
    return params, specialization, options
c                Z    i | ](\  }}|j         t          j        j        ud | |j         )S )default_r  )r   r4   r   s      rD   r  z2create_function_from_signature.<locals>.<dictcomp>  sE       D%= 1 777 	45=777rE   specialize_implbackendrX   dynamic_func)r'   
parametersziprY   r   appendr   r   r   r   r/   r^   joinr*   r   r  r   rX   exec)sigkparamsr&  specializationr4   kpr   r   r   rJ   r   	func_bodyfunc_namespacer%  s                 rD   create_function_from_signaturer3    s    s~#g,,....N++--w77 0 0b? 	0!!":4":":":;;;;!#9vv'H$&$8DfJ!@LGGfEXdXXhXX*XXPUXXXC! 0b0#66 +)T11R5G5K|5[5[%*
 M"))*Rr/A*R*R*R*R*RSSSS #))*Kr/A*K*K*KLLLL%%h//// b
aC))DS#.*>*>*@*@!A!ABBk]RSS 		QQ3>;N;N;P;PQQQRR  xx//  I >//11  N -O(7N$% 'N9$/N=! 	N### .))rE   c                $    | j          d| j         S )N.)rL   r   fns    rD   get_full_namer8    s    m//bo///rE   c                      e Zd Zd Zd Zedd            Zd Zd Zed             Z	d	 Z
d
 Zd Z eee          ZdS )rX   c                   || _         t          j        |          | _        	 t          j        |          \  | _        | _        n"# t          $ r}t          d          |d }~ww xY wt          |          | _	        t          j                    | _        t          j        d                    | j                            }|t!          j        d|t           j                                                  d          }|| _        d | _        i | _        |j        | _        |j        | _        |j        | _        |j        | _        |j        | _        d S )Nz1@jit functions should be defined in a Python filerM   z^def\s+\w+\s*\()r7  rN   r  getsourcelinesraw_srcstarting_line_numberOSError
ValueErrorr8  _fn_name	threadingRLock
_hash_locktextwrapdedentr+  research	MULTILINEstart_srchashr?   r   r[   r   __globals__rL   )rA   r7  erB   s       rD   r3   zJITCallable.__init__  s,    *2..	Y6=6LR6P6P3DL$33 	Y 	Y 	YPQQWXX	Y%b))#/++ obggdl3344").R\BBHHJJKKL		 TV zO>-s   !A 
A#AA#c                N    | j         t          j        | j                  j        z  S rG   )rL  rN   getclosurevarsr7  r:   rI   s    rD   get_capture_scopezJITCallable.get_capture_scope  s     '"8"A"A"KKKrE   r"   r^   c                4   | j         5  | j        | j        cd d d            S d| j         | _        t          j        | j                  j        }t          | j        | j        || j	                  }|
                    |                                            |j        t          | j                  z   | _        t          t!          |j                                                            | _        ddlm | xj        t          fd| j                                        D                       z  c_        t+          j        | j                            d                                                    | _        d d d            n# 1 swxY w Y   | j        S )Nz
recursion:)r4   r9   r:   rB   r   rg   c                J    g | ]\  \  }}\  }}t          |          ||f S r   )r/   )r   r4   ra   rr   rh   s       rD   r   z)JITCallable.cache_key.<locals>.<listcomp>  sH     = = ="5)4Xc1!+C!;!;=tSk = = =rE   r%   )rC  rK  r@  rN   rO  r7  r:   r!   rL  rB   r   parserJ   r^   r=  dictsortedr?   r  r   rh   r5   r6   r7   rH   )rA   r:   dependencies_finderrh   s      @rD   r]   zJITCallable.cache_key  s    _ 	N 	Ny$y	N 	N 	N 	N 	N 	N 	N 	N
 5T]44DI.tw77AI"4$-QUQamv9=#C #C #C%%djjll333+/#d6O2P2PPDI$(0C0T0Z0Z0\0\)])]$^$^D!666666II = = = =9=9N9T9T9V9V= = = > > >II  ty'7'7'@'@AAKKMMDI#	N 	N 	N 	N 	N 	N 	N 	N 	N 	N 	N 	N 	N 	N 	N$ ys   FEFFFc                *    t          | j                  S rG   )rK  r]   rI   s    rD   __hash__zJITCallable.__hash__  s    DN###rE   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   )r{   rS  rJ  r/   Moduler'   bodyFunctionDef)rA   trees     rD   rS  zJITCallable.parse  sg    y##$
+++++49~~""""$)A,88888rE   c                $    ddl m}  ||           S )Nr   )constexpr_type)r   r_  )rA   r_  s     rD   rl   zJITCallable.type  s$    777777~d###rE   c                "    d| _         || _        dS )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)rK  rJ  )rA   new_srcs     rD   _unsafe_update_srczJITCallable._unsafe_update_src  s     				rE   c                     t          d          )NzqCannot set attribute 'src' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)AttributeErrorrI   s    rD   _set_srczJITCallable._set_src!  s     ( ) ) 	)rE   c                    | j         S rG   )rJ  rI   s    rD   _get_srczJITCallable._get_src&  s
    yrE   )fgetfsetNr   )r[   rL   r   r3   rP  r   r]   rX  rS  rl   rb  re  rg  rB   r   rE   rD   rX   rX     s         (  (  (DL L L    X,$ $ $   $ $ X$  ) ) )
   (x
0
0
0CCCrE   rX   c                  .    e Zd ZU ded<   ded<   ded<   dS )JitFunctionInfor   rS   r^   r4   JITFunctionjit_functionN)r[   rL   r   r   r   rE   rD   rk  rk  ,  s3         IIIrE   rk  c                    t          |          t          |          f}|                     |d           }||S fdt           |                    t          |          z   }|| |<   |S )Nc                4   t          | t                    rfd| D             S t          |           rfd| D             } | j        | S t          | t                    rt	          fd| D                       S t          | t
                    r| j        S | S )Nc                &    g | ]} |          S r   r   r   r   replace_callabless     rD   r   z@compute_cache_key.<locals>.replace_callables.<locals>.<listcomp><  s%    :::s%%c**:::rE   c                &    g | ]} |          S r   r   rq  s     rD   r   z@compute_cache_key.<locals>.replace_callables.<locals>.<listcomp>>  s%    ===#((--===rE   c              3  .   K   | ]} |          V  d S rG   r   rq  s     rD   	<genexpr>z?compute_cache_key.<locals>.replace_callables.<locals>.<genexpr>A  s/      ??C**3//??????rE   )r/   r*   r   rC   tuplerX   r]   )r  resultsrr  s     rD   rr  z,compute_cache_key.<locals>.replace_callables:  s    c4   	!::::c::::3 	!=======G 3='**U## 	!????3??????[)) 	!= 
rE   )rv  r^   rx   )kernel_key_cacher/  r  r   r]   rr  s        @rD   compute_cache_keyry  3  s      #g,,
/C $$S$//I
 
 
 
 
 %%n5566WEI%SrE   c                    t          | t                    s| S t          |           D ]\  }}t          |          | |<   t	          |           S rG   )r/   r*   	enumerateconvert_to_tuple_if_listrv  )iteminested_values      rD   r|  r|  K  sS    dD!!  %T?? 9 9<*<88Q;;rE   c                  h     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 )rl  c                    dS )NFr   rI   s    rD   is_gluonzJITFunction.is_gluonY  s    urE   r"   bool | Nonec	                   |sd S | j         j        }	| j         j        }
d                    d t	          | j        |d                   D                       }|	 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 )Nr  c                ,    g | ]\  }}|j          d | S )z: r4   )r   r   r   s      rD   r   z*JITFunction._call_hook.<locals>.<listcomp>l  s,    ___%*4444___rE   r   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   reprr7  compileis_manual_warmupalready_compiled)r7  r   rL   r+  r)  paramsr  r  r  r  r  r8  r  r  rk  )rA   hookr   r  r  r  r  r  r  r4   rS   	arg_reprsr  	full_namer  r   s                   rD   
_call_hookzJITFunction._call_hook\  s     	4w##II__c$+WZ[\W]F^F^___``	  k  k7#4  k  kAQ  k  k`g`r  k  k  HO  H`  k  k  |C  |[  k  k  _h  k  k  k!$'**	;IyR[]def]gipruvv #" *(!, ' 8'.'F".#6"
 
 tvtT22C*6*&"
 
 
 	
rE   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)ro   pre_run_hooksr*  )rA   r  s     rD   add_pre_run_hookzJITFunction.add_pre_run_hook  s3    
 ~~!!$'''''rE   c                    ddl m}m}m}m} t
          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_targetr3  r  r  )rA   r  r  r  r  r   r&  binders           rD   create_binderzJITFunction.create_binder  s     	POOOOOOOOOOO1133,v&&,"/WUU2vw..rE   c                                        |          }d | j        D             }d |D             }d t          ||          D             }d|vs
J d            d|vs
J d            d|vs
J d	            |D ]!}	|	|j        vr|	|vrt	          d
|	z            "t          |d           }
fd|
D             }
d |D             t          d           }fd|D             }|||
|fS )Nc                    g | ]	}|j         
S r   r  r  s     rD   r   z*JITFunction._pack_args.<locals>.<listcomp>  s    ///a16///rE   c                    g | ]
}|d          S )r   r   r  s     rD   r   z*JITFunction._pack_args.<locals>.<listcomp>  s    000A1Q4000rE   c                    i | ]\  }}||	S r   r   )r   r_   vs      rD   r  z*JITFunction._pack_args.<locals>.<dictcomp>  s    >>>fq!Q>>>rE   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   )ra   rr   s     rD   r   z(JITFunction._pack_args.<locals>.<lambda>  s    3+;M rE   c           	     p    i | ]2}|t          t                                                    |          3S r   )r   r*   r   )r   path
bound_argss     rD   r  z*JITFunction._pack_args.<locals>.<dictcomp>  s;    fffSWd-d:3D3D3F3F.G.GNNfffrE   c                    g | ]
}|d          S )r   r   r  s     rD   r   z*JITFunction._pack_args.<locals>.<listcomp>  s    111QAaD111rE   c                ,    t          |t                    S rG   )r/   r^   )ra   r  s     rD   r   z(JITFunction._pack_args.<locals>.<lambda>  s    Z35G5G rE   c           	     X    i | ]&}|                     t          |                    'S r   )
parse_attrr   )r   r_   attrvalsr&  s     rD   r  z*JITFunction._pack_args.<locals>.<dictcomp>  s4    VVV1G&&'81'E'EFFVVVrE   )parse_optionsr  r)  r  KeyErrorr   )rA   r&  r   r  r/  r  sigkeyssigvalsr  r_   
constexprsr  r  s    ` `        @rD   
_pack_argszJITFunction._pack_args  si   ''////4;///00000>>GW(=(=>>>	F***,k***v%%%'a%%%v%%%'a%%% 	Y 	YA(((Qg-=-=SVWWXXX"7,M,MNN
ffff[efff
11.111h(G(GHHVVVVVPUVVV	:u44rE   c                  |                     d| j                  pt          j        j        |d<   t          j        j        |d<   t          j                                        }t          j        	                    |          }| j
        D ]
} ||i | | j        |         \  }}	}
}} ||i |\  }}}t          |	||          }|                     |d           }|=|                     |||||          \  }}}}|                     |||||||          }|d S t                      }| j                                        D ]?\  \  }}\  }}|                     ||          x}|k    rt%          d| d| d|           @|s|J t'          |          r ||          }t)          |          }|d         }|dk    r|d         nd}|dk    r|d         nd}t+          |d	          r|                                } |j        ||g|                                R  } |j        |||||j        |j        |t          j        j        t          j        j        g	|                                R   |S )
Ndebuginstrumentation_moderV   z1 has changed since we compiled this kernel, from z to r   r   r   result)rx   r  r   runtimecompilationr  r   r  get_current_deviceget_current_streamr  device_cachesry  r  _do_compileobjectr?   r  rZ   ro   r'   hasattrr  launch_metadatar   r   functionpacked_metadatalaunch_enter_hooklaunch_exit_hook) rA   r   r   r   r   r  r  r  kernel_cacherx  r   r&  r  r  r/  r  r   kernelr  r  r  not_presentr4   ra   rr   globals_dictnewVal	grid_sizegrid_0grid_1grid_2r  s                                    rD   r   zJITFunction.run  s    **Wdj99PU]=Pw).):)O%& 113311&99 & 	" 	"DD$!&!!!!BFBTU[B\?& /5fd.Ef.E.E+
NG 0.'JJ!!#t,, >48OOGVU_aoDK5M 5M1GY
E %%c9fj'SXZ`aaF~t 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  	n###~~ (tJ''D		I!WF )AT!WW1F )AT!WW1Fvx(( )4f4T6XJDUDUDWDWXXXOFJvvvvvH^`o}68VnYcYjYjYlYln n n nrE   c                H    | j         | j        n|                      |          S rG   )_reprr@  )rA   ra   s     rD   r  zJITFunction.repr  s     $
 2t}}

1ErE   Nc	           	     z   |r|ng }|r|ng }t                                          |           |j        | _        || _        || _        || _        || _        || _        g | _	        t          | j        j                                                  D ]I\  }	}
|	|v p|
j        |v }|	|v p|
j        |v }| j	                            t!          |	|
||                     Jt#          | j                  | _        d | _        || _        || _        d | j	        D             | _        d | j	        D             | _        g | _        d S )Nc                    g | ]	}|j         
S r   r  r   ps     rD   r   z(JITFunction.__init__.<locals>.<listcomp>  s    666Q!&666rE   c                *    g | ]}|j         	|j        S r   )r   r   r  s     rD   r   z(JITFunction.__init__.<locals>.<listcomp>  s!    HHHQH15HHHrE   )r2   r3   rL   rS   versionr   r   r  r  r  r{  r  r(  r   r4   r*  r   r   r  r  r  r  rW   	arg_namesr  r  )rA   r7  r  r   r   r  rW   r  r  r~  r   dnsdns_oarC   s                rD   r3   zJITFunction.__init__  se   1BJ--Ki)q)G)Goq&m!2.L+
.!$.";"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6$+666HH$+HHH  rE   c           	        dd l }dd lm t          j                                        }|                    |          }|d         | j        k    r t          d|d          d| j                   t          t          |d                   }|d         }fdt          ||          D             }t          t          |d                   }|d	         }	t          t          ||	                    }
d
 |d                                         D             }d |d                                         D             }|d         }| j        |         \  }}}}}|                    |          }|                     ||||||
d          S )Nr   r4   zSpecialization data is for z but trying to preload for r  r	  c                    i | ]k\  }}|j                             |          r                     |          n5t          |t                    rd |v r                    |d                    n|lS rg   )r   is_dtyper/   rT  rh   )r   r   r   tls      rD   r  z'JITFunction.preload.<locals>.<dictcomp>  s     
 
 
 U !x0077 lBHHUOOO0:5$0G0GkK[`L`L`BLL{+,,,fk
 
 
rE   r
  r  c                4    i | ]\  }}|t          |          S r   )r|  r   s      rD   r  z'JITFunction.preload.<locals>.<dictcomp>(  s'    rrrjc5S2599rrrrE   r  c                b    i | ],\  }}|t          |t                    rt          |          n|-S r   )r/   r*   rv  r   s      rD   r  z'JITFunction.preload.<locals>.<dictcomp>)  sG     
 
 
U E4!8!8Cue
 
 
rE   r  r   T)r   )r  triton.languager   r   r  r  loadsr@  rZ   r   rv  r)  rT  r  r  r  r  )rA   r  r  r  deserialized_objr  r	  r  r
  r  r  r  r  r   ra   r&  r  s                   @rD   preloadzJITFunction.preload  s   $$$$$$1133::&9::F#t}44r.>v.Frrcgcprrt t tE#3O#DEE(9
 
 
 
 "-??	
 
 

  0 >??
%l3
SZ0011 srL\]hLiLoLoLqLqrrr	
 
.y9??AA
 
 
 u%"081a!''00   
 
 	
rE   c           
     H     j                  \  }}	}                     t          j        j        g          rd S                                 t          j                                        }
|
Nt                      t          |	          } fd} f	d}|
                    |||          }nN                     j                  }|<                        t          j        j        g           |S )Nc                 @                         j                   S )N)r   r  	_env_vars)r  r  )env_varsr  rA   rB   r   s   rD   async_compilez.JITFunction._do_compile.<locals>.async_compileG  s!    ||C@P\d|eeerE   c           
     j   	 | <                        t          j        j        g	           d S rG   )r  r   r  jit_post_compile_hook)
r  r  r  r  r  r   r  rA   r  r   s
    rD   finalize_compilez1JITFunction._do_compile.<locals>.finalize_compileJ  sH    $*S! CS)U[]gip!&1 1 1 1 1rE   )r   r  )r  r  r   r  jit_cache_hookr  r   active_moderx   r   r   submitr  r  r  )rA   r   r  r  r  r  r  r   ra   r&  
async_moder]   r  r  r  r  r  rB   r   s   ````````       @@@@rD   r  zJITFunction._do_compile:  s   .2.@.H+a!??5=7iQ[]dglfmouvv 	4nnT9j%@@#/3355
!688H%c7GXFFIf f f f f f f f f1 1 1 1 1 1 1 1 1 1 1 1 1
  &&y-AQRRFF\\#fg>N\OOF &LOOEM?iQWYcelotnu"$ $ $rE   c                     t          d          )Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rZ   rA   r   r   s      rD   __call__zJITFunction.__call__W  s    WXXXrE   c                2    d| j          d| j        j         dS )NzJITFunction(:r  )rS   r7  r   rI   s    rD   __repr__zJITFunction.__repr__Z  s"    CdkCCDG,@CCCCrE   )r"   r  )NNNNNNN)r[   rL   r   r  r  r  r  r  r   r  r3   r  r  r  r  r   r   s   @rD   rl  rl  W  s         ,
 ,
 ,
 ,
\( ( (/ / /5 5 503 3 3jF F F mq;?"  "  "  "  "  " H%
 %
 %
N  :Y Y YD D D D D D DrE   rl  r7  JITFunction[T]c                    d S rG   r   r6  s    rD   jitr   c  s    CrE   r  r  r  r   r   r  rW   r  Optional[Callable]r  r   Optional[Iterable[int | str]]r   r  Optional[bool]rW   Callable[[T], JITFunction[T]]c                    d S rG   r   r  s          rD   r   r   h  s	     CrE   Optional[T]KernelInterface[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
    r7  r   r"   r  c           
         t          |           sJ t          j        j        rddlm}  ||           S t          |           S )Nr   )InterpretedFunction)r  r   r   r  rW   r  r  )ro   r   r  	interpretinterpreterr  rl  )	r7  r  r  r   r   r  rW   r  r  s	     rD   	decoratorzjit.<locals>.decorator  s    ||=" 	888888&&r7N_Fdlq08tUdf f f f "3/M! /	 	 	 	rE   Nr7  r   r"   r  r   )	r7  r  r  r  r   r   r  rW   r  s	    ``````` rD   r   r   v  sb    :           & 
~y}} rE   c                  b    e Zd ZdZed             Zd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)rC   r[   rL   r   r   s    rD   r   zMockTensor.wrap_dtype  s/    =!W,,71J1Jc??"
rE   Nc                ,    |dg}|| _         || _        d S )Nr   )r   shape)rA   r   r  s      rD   r3   zMockTensor.__init__  s     =CE



rE   c                    dg}| j         dd          D ] }|                    |d         |z             !t          t          |                    S )Nr   r   )r  r*  rv  reversed)rA   stridessizes      rD   stridezMockTensor.stride  sV    #JqrrN 	/ 	/DNN72;-....Xg&&'''rE   c                     dS Nr   r   r   rE   rD   data_ptrzMockTensor.data_ptr      qrE   c                     dS r  r   r   rE   rD   	ptr_rangezMockTensor.ptr_range  r  rE   rG   )
r[   rL   r   r   staticmethodr   r3   r  r  r  r   rE   rD   r   r     s         
   \
   ( ( (   \   \  rE   r   c                  L    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dS )TensorWrapperc                t    || _         || _        |j        | _        |j        | _        | j        j        | _        d S rG   )r   basedatar  r  )rA   r$  r   s      rD   r3   zTensorWrapper.__init__  s1    
	I	kY_


rE   c                4    | j                                         S rG   )r$  r  rI   s    rD   r  zTensorWrapper.data_ptr  s    y!!###rE   c                      | j         j        | S rG   )r$  r  )rA   r   s     rD   r  zTensorWrapper.stride  s    ty&&rE   r"   r^   c                (    d| j          d| j         dS )NzTensorWrapper[r  r  )r   r$  rI   s    rD   __str__zTensorWrapper.__str__  s    :
::di::::rE   c                4    | j                                         S rG   )r$  element_sizerI   s    rD   r+  zTensorWrapper.element_size  s    y%%'''rE   c                Z    t          | j                                        | j                  S rG   )r"  r$  cpur   rI   s    rD   r-  zTensorWrapper.cpu  s    TY]]__dj999rE   c                D    | j                             |j                    d S rG   )r$  copy_)rA   others     rD   r/  zTensorWrapper.copy_  s    	
#####rE   c                Z    t          | j                                        | j                  S rG   )r"  r$  cloner   rI   s    rD   r2  zTensorWrapper.clone  s     TY__..
;;;rE   c                \    t          | j                            |          | j                  S rG   )r"  r$  tor   )rA   r  s     rD   r4  zTensorWrapper.to  s"    TY\\&114:>>>rE   c                \    t          | j                            |          | j                  S rG   )r"  r$  	new_emptyr   )rA   sizess     rD   r6  zTensorWrapper.new_empty  s$    TY0077DDDrE   Nr   )r[   rL   r   r3   r  r  r)  r+  r-  r/  r2  r4  r6  r   rE   rD   r"  r"    s        % % %$ $ $' ' '; ; ; ;( ( (: : :$ $ $< < <? ? ?E E E E ErE   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 r5  )r/   r"  r$  r   r  r   rl   )tensorr   s     rD   reinterpretr:    s    &-(( AFK%%%; !e444		$	$ AVU+++?V???@@@rE   c                0   | }t          |t                    s|j        }t          |t                    |j        j        j        }|j        }t          |j                  D ]3\  }}|                                	                    d          r||z  } n4||fS )Nzdef )
r/   rX   r7  __code__co_filenamer=  r{  r<  r   rQ   )r7  base_fn	file_name
begin_lineidxlines         rD   get_jit_fn_file_linerC    s    G+.. * +.. 
#/I-J w//  	T::<<""6** 	#JE	 j  rE   c                  0    e Zd Zd Zed             Zd ZdS )BoundConstexprFunctionc                "    || _         || _        d S rG   )__self____func__)rA   instancer7  s      rD   r3   zBoundConstexprFunction.__init__  s     rE   c                    | j         j        S rG   )rH  r]   rI   s    rD   r]   z BoundConstexprFunction.cache_key  s    }&&rE   c                .     | j         | j        g|R i |S rG   )rH  rG  r  s      rD   r  zBoundConstexprFunction.__call__#  s&    t}T]<T<<<V<<<rE   N)r[   rL   r   r3   r   r]   r  r   rE   rD   rE  rE    sM           ' ' X'= = = = =rE   rE  c                  0     e Zd Z fdZd ZdddZ xZS )ConstexprFunctionc                J    t                                          |           d S rG   )r2   r3   )rA   r7  rC   s     rD   r3   zConstexprFunction.__init__)  s!    rE   c                *    |t          ||           S | S rG   )rE  )rA   r  objclasss      rD   __get__zConstexprFunction.__get__,  s    ?)#t444rE   N)	_semanticc                   ddl mm} fd|D             }fd|                                D             } | j        |i |}||S t
          j        j        r|S  ||          S )Nr   )_unwrap_if_constexprrh   c                &    g | ]} |          S r   r   )r   r  rT  s     rD   r   z.ConstexprFunction.__call__.<locals>.<listcomp>5  s%    666A$$Q''666rE   c                .    i | ]\  }}| |          S r   r   )r   r_   r  rT  s      rD   r  z.ConstexprFunction.__call__.<locals>.<dictcomp>6  s+    JJJ!Q!))!,,JJJrE   )r   rT  rh   r  r7  r   r  r  )rA   rR  r   r   rh   resrT  s         @rD   r  zConstexprFunction.__call__2  s    HHHHHHHH6666666JJJJ6<<>>JJJ dgt&v&&J =" 	Jy~~rE   )r[   rL   r   r3   rQ  r  r   r   s   @rD   rM  rM  '  se               )-         rE   rM  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.
    )rM  r6  s    rD   constexpr_functionrY  E  s     R   rE   r   )Fr  )r  r  r  r  r   r  r   r  r  r  rW   r  r"   r  rG   )r7  r  r  r  r  r  r   r  r   r  r  r  rW   r  r"   r  )F
__future__r   r   r{   r0   r5   rN   r   rA  rF  rD  collectionsr   dataclassesr   	functoolsr   typingr   r	   r
   r   r   r   r   r   r   triton.backendsr   typesr   rM   r   r   r   _utilsr   r   r   r   cacher   triton._C.libtritonr   r   r=   r<   r   NodeVisitorr!   r   r   r   r   r  r3  r8  rX   rk  ry  r|  rl  r   r   r"  r:  rC  rE  rM  rY  r   rE   rD   <module>re     s   , , , , , , , , 



            				  # # # # # # ! ! ! ! ! ! % % % % % % ] ] ] ] ] ] ] ] ] ] ] ] ] ] ] ] ] ] ] ] ] ] ' ' ' ' ' '                         ` ` ` ` ` ` ` ` ` ` ` `             W W W W W W W W!3GCLLg! g! g! g! g! g! g! g!^D D D D4/> /> /> /> /> /> /> />dT T T TZ Z Z Z Zgaj Z Z Z&  "9* 9* 9*x0 0 0b1 b1 b1 b1 b1 b1 b1 b1J          0	 	 	DD DD DD DD DD+q1 DD DD DDX 
   
 
 #*.7;DH #
 
 
 
 
 

 4 #*.7;DH #4 4 4 4 4 4x       B"E "E "E "E "E "E "E "EJA A A! ! !$= = = = =[ = = =       <! ! ! ! !rE   