
    rh              	         d dl mZmZ d dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
mZ d dlmZ d dlmZ d dlmZmZmZmZmZmZmZmZmZmZ d dlmZ d dlmZ d	d
lm Z  d	dl!m"Z" d	dl#m$Z$m%Z%m&Z&m'Z' e(d e)d        Z* ed      Z+ G d dejX                        Z-d)dZ. G d d      Z/i Z0g Z1d Z2d*dZ3 G d dee+         Z4d Z5d Z6d Z7e G d d             Z8 G d de4e+         Z9ed+d       Z:edddddddd 	 	 	 	 	 	 	 	 	 	 	 	 	 d,d!       Z:	 d-dddddddd 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 d.d"Z: G d# d$      Z; G d% d&      Z<d' Z=d( Z>y)/    )annotationsdivisionN)defaultdict)	dataclass)cached_property)
CallableGenericIterableOptionalTypeVarUnionoverloadDictAnyTupleTensorDescriptor)
ModuleType   )knobs)driver)find_paths_ifget_iterable_pathtype_canonicalisation_dictcanonicalize_dtypez.runtime.jitTc                  v     e Zd ZdZd fdZed        Zd Zd Zd Z	d Z
d Zd	 Zd
 Zd Zd Zd Zd Z xZS )DependenciesFindera  
    This AST visitor is used to find dependencies of a JITFunction. This can
    be used to invalidate a JITFunction's hash when its source code -- or
    that of its dependencies -- changes.

    This visitor also keeps track of the global variables touched by the
    JITFunction.  When we launch the kernel, we check that these have the same
    values as they did when we ran this visitor.  If not, we raise an error (or
    otherwise we could recompile).
    c                    t         |           || _        t        j                  |j                  d            | _        || _        || _        h d| _	        i | _
        d| _        y )Nutf-8>
   intlenmaxminlistfloatprintrangegetattr
isinstanceF)super__init__namehashlibsha256encodehasherglobals	nonlocalssupported_python_builtinsused_global_valsvisiting_arg_default_value)selfr-   r2   r3   src	__class__s        e/var/www/html/ai-insurance-compliance-backend/venv/lib/python3.12/site-packages/triton/runtime/jit.pyr,   zDependenciesFinder.__init__)   s\    	nnSZZ%89 "*
&. TV*/'    c                6    | j                   j                         S N)r1   	hexdigestr7   s    r:   retzDependenciesFinder.retN   s    {{$$&&r;   c                    t        j                  |j                        ryt        |dd      }|j	                  t
              S )NT
__module__ )inspect	isbuiltinfuncr)   
startswithTRITON_MODULE)r7   noderF   modules       r:   _is_triton_builtinz%DependenciesFinder._is_triton_builtinR   s6    TYY'|R0  //r;   c                >   t        |t              r| j                  j                         |j                  j                         z  D ]_  }|\  }}| j                  |   \  }}|j                  |   \  }}||k7  s2t	        d| d| d| j
                   d|j                   d| d       | j                  j                  |j                         |j                  }|t        t        |dd            z  }| j                  j                  |j                  d	             y y )
NGlobal variable z has value z when compiling z, but inner kernel z has conflicting value z7 from when it was first compiled.  This is not allowed.noinlineFr    )r*   JITFunctionr5   keysRuntimeErrorr-   __name__update	cache_keystrr)   r1   r0   )r7   rF   kvar_name_v1v2func_keys           r:   _update_hashzDependenciesFinder._update_hashX   s3   dK( **//1D4I4I4N4N4PP !--a0A--a0A8&*8*KtCSTXT]T]S^^qrvrr  rA  AX  Y[  X\  \S  T  !!(()>)>?~~HGD*e<==HKKxw78 )r;   c                    t        |j                        t        j                  u r|j                  S |j                   j
                  v ry  fd} ||j                        \  }}| j                  st        |      t        urnt        |t              s^t        |dd      sQ|j                   j                  vr9t        j                  |      |f j                  |j                  t	        |      f<    j                  |       |S )Nc                    j                   j                  | d       }||j                   fS j                  j                  | d       }||j                  fS y)N)NN)r2   getr3   )r-   valr7   s     r:   name_lookupz2DependenciesFinder.visit_Name.<locals>.name_lookupr   sZ    ,,""4.CDLL((..$$T40CDNN**r;   __triton_builtin__F)typectxastStoreidlocal_namesr6   r   r*   rO   r)   r4   copyr5   r\   )r7   rI   ra   r`   var_dicts   `    r:   
visit_NamezDependenciesFinder.visit_Namej   s    >SYY&77N77d&&&	 $DGG,X
 O 77 IZ/ #34WSJ^`e=fGG4#A#AA>Biinh=WD!!477BxL"9:#
r;   c                ^    |j                   D cg c]  }| j                  |       c}S c c}w r=   )eltsvisit)r7   rI   elts      r:   visit_TuplezDependenciesFinder.visit_Tuple   s$     ,0995C

3555s   *c                X   | j                  |j                        }t        |t        j                        r6| j                  |j                        }t        |t        j                        r6|t        |dd      t        k(  ry t        ||j                        }| j                  |       |S )NrR   rC   )	rn   valuer*   re   	Attributer)   rH   attrr\   )r7   rI   lhsr@   s       r:   visit_Attributez"DependenciesFinder.visit_Attribute   s    jj$cmm,**SYY'C cmm,;73
B7=Hc499%#
r;   c                    |j                   j                   D ch c]  }|j                   c}| _        | j                  |       y c c}w r=   )argsargrh   generic_visit)r7   rI   ry   s      r:   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)r6   rn   )defaultsexprr7   s     r:   visit_defaultsz:DependenciesFinder.visit_arguments.<locals>.visit_defaults   sS    8::::26/$ )D'

4() 38/%/s   < < 	A)
	itertoolschainposonlyargsrx   vararg
kwonlyargsrn   kw_defaultskwargr~   )r7   rI   r   ry   s   `   r:   visit_argumentsz"DependenciesFinder.visit_arguments   s    	8 ??4#3#3TYYQUQ\Q\bdfjfufuv 	CJJsO	 	t''(::!JJtzz"t}}%r;   c                    | j                  |      }t        |t              r| xj                  t	        |      z  c_        y | j                  j                  |       y r=   )rn   r*   r%   rh   setadd)r7   rI   targets      r:   visitAssnTargetz"DependenciesFinder.visitAssnTarget   sE     D!fd#F+  (r;   c                    t        |j                        dk7  rt        d      | j                  |j                  d          | j	                  |       y )N   z2Simultaneous multiple assignment is not supported.r   )r"   targets	TypeErrorr   rz   r7   rI   s     r:   visit_AssignzDependenciesFinder.visit_Assign   sG    t||!
 PQQT\\!_- 	4 r;   c                \    | j                  |j                         | j                  |       y r=   r   r   rz   r   s     r:   visit_AnnAssignz"DependenciesFinder.visit_AnnAssign   $    T[[) 	4 r;   c                \    | j                  |j                         | j                  |       y r=   r   r   s     r:   	visit_ForzDependenciesFinder.visit_For   r   r;   )returnNone)rR   rB   __qualname____doc__r,   propertyr@   rK   r\   rk   rp   rv   r{   r   r   r   r   r   __classcell__r9   s   @r:   r   r      s[    	#0J ' '09$%N6
!
&@)!!!r;   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._trC   )triton.language.corelanguagecorer*   rU   striprG   removeprefix_normalize_tyendswithpointer_type
element_tydtyper-   rc   rR   r   r_   replace)tyr   s     r:   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Cr;   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 r=   )num_paramdo_not_specializedo_not_specialize_on_alignment)r7   r   paramr   r   s        r:   r,   zKernelParam.__init__
  s     !2.L+r;   c                .    | j                   j                  S r=   )r   r-   r?   s    r:   r-   zKernelParam.name  s    {{r;   c                    | j                   j                  r1| j                   j                  t        j                  j                  k(  ryt        | j                   j                        S )NrC   )r   
annotationrD   	Parameteremptyr   r?   s    r:   r   zKernelParam.annotation  sD    {{%%)?)?7CTCTCZCZ)ZT[[3344r;   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   rC   )r   rG   r   r   values)r7   as     r:   annotation_typezKernelParam.annotation_type  s]    OO<<!"A\\#!"A.55788??"r;   c                    d| j                   v S N	constexpr)r   r?   s    r:   is_constexprzKernelParam.is_constexpr&  s    doo--r;   c                r    | j                   ryd| j                  v xs | j                  j                  d      S )NFr   r   )r   r   rG   r?   s    r:   is_constzKernelParam.is_const*  s1    $//)MT__-G-G-MMr;   c                .    | j                   j                  S r=   )r   defaultr?   s    r:   r   zKernelParam.default0  s    {{"""r;   c                d    | j                   j                  t        j                  j                  k7  S r=   )r   r   rD   r   r   r?   s    r:   has_defaultzKernelParam.has_default4  s#    {{""g&7&7&=&===r;   N)r   r!   r   zinspect.Parameterr   boolr   r   r   rU   )rR   rB   r   r   r,   r   r-   r   r   r   r   r   r   r    r;   r:   r   r     s    LM15M     5 5
   . . N N
 # # > >r;   r   c                8     ddl m ddlm d fd	S )Nr   r   r   r   c                     yt         t              ryt         t              r7|r  d|      nd } dk(  r|ryd k  r	 dk  rd	|fS d
 k  r	 dk  rd|fS d|fS t         t              ryt	         d      rZ j
                  |f}t        j                  |d       }|!|d   rdndt        |d         z   }|t        |<   |r  d|      nd }||fS t         t              rd j                  fS t               rd fS t	         d      ryt         t              rP D cg c]
  } |       }} fd}	 |	|D cg c]  }|d   	 c}      }
 |	|D cg c]  }|d   	 c}      }|
|fS t         t              rTt	         j                  d      sJ t         j                  j
                        }d| t         j                         dd fS t               rat	         j                  d      sJ t         j                  j
                        }d| t         j                         d j                   dd fS t#        dt%               z        c c}w c c}w c c}w )N)r   N)u1Nr!   )alignr   )r   r   i   ii32l            l    u64i64)fp32Ndata_ptrr   r   r   tensorr   tma_desc_cpu_ptr)	nvTmaDescNc                N    t        d      r t              |  S t        |       S )N_fields)hasattrrc   tuple)valsry   s    r:   <lambda>zAcreate_specialize_impl.<locals>.specialize_impl.<locals>.<lambda>d  s&    '#y:Qid3i&6 W\]aWb r;   ztensordesc<>,zUnsupported type: %s)r*   r   r!   r&   r   r   	dtype2strr_   r   rO   rT   r   r   baser%   block_shapelayoutr   rc   )ry   r   specialize_valuer   keydskresxspec
make_tupletysrP   innerGluonTensorDescriptorr   specialize_extraspecialize_impls   `            r:   r   z/create_specialize_impl.<locals>.specialize_implB  sj   ;&T"S!?O"3U;UYCax,'SSI%5s|###"2s|#s|#U#!S*%99h'C--T*C{"1vt32DSV2LL!$	#BR"3>X\C:[)//Y'%%S,-&U#0341OA&4D4bJD1qad12CT2qt23D;-.388Z000&sxx~~6E!%coo)>(?qA4HH23388Z000&sxx~~6E!%coo)>(?qaPRVWW2T#Y>?? 512s   $I II")FTT)r   r   'triton.experimental.gluon.nvidia.hopperr   )r   r   r   r   s   `@@@r:   create_specialize_implr   =  s    $a/@ /@b r;   c                    t        t              dk(  rt        j                  t        d              t        d   } || |      d   S )Nr   c                     y r=   r   )rX   kwargss     r:   r   zmangle_type.<locals>.<lambda>x  s    r;   )r   )r"   specialize_impl_cacheappendr   )ry   
specializer   s      r:   mangle_typer   v  sA    
 !Q&$$%;<T%UV+A.O3<Q??r;   c                       e Zd ZU ded<   ddZy)KernelInterfacer   runc                      fdS )z
        A JIT function is launched with: fn[grid](*args, **kwargs).
        Hence JITFunction.__getitem__ returns a callable proxy that
        memorizes the grid.
        c                 .     j                   | dd|S )NFgridwarmup)r  )rx   r   r  r7   s     r:   r   z-KernelInterface.__getitem__.<locals>.<lambda>  s    xtxx$T%'YRX'Y r;   r   )r7   r  s   ``r:   __getitem__zKernelInterface.__getitem__  s     ZYr;   N)r   r   )rR   rB   r   __annotations__r  r   r;   r:   r   r   }  s    	
FZr;   r   c           
        |j                         D ci c],  \  }}||j                  j                  dk(  rt        |      n|. }}}dd l}| ||j                         D cg c]  }t        |       c}t        |j                               |j                         D cg c]  }t        |       c}t        |j                               |j                  d}	|j                  |	      }
|
S c c}}w c c}w c c}w )Nr   r   )r-   	signatureconstant_keysconstant_vals
attrs_keys
attrs_valsoptionsr   )
itemsr9   rR   rU   jsonrP   r%   r   __dict__dumps)r-   r
  	constantsattrsr  r   rr   r  r   objserialized_objs              r:   serialize_specialization_datar    s    enetetevwWaWZ\aEOO$<$<$Gc%jURwIw9QZQ_Q_Qa?bAQ?bY %**,0OQa0O_cdidpdpdr_s##CC
 ZZ_N x @c0Os   1C, C2C7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<   t+        |j,                        |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(, r   Nr   )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   rD   r   r   )r   s    r:   r   z0create_function_from_signature.<locals>.<lambda>  sA    AaDLLG,=,=,C,CCAaD AaD6QZ[\]^[_Z`Ia r;   z
def dynamic_func(z	**optionsz):
    params = {'z': z}
    specialization = [r   z-]
    return params, specialization, options
default_rO   r   dynamic_func)r"   
parametersziprP   r   r   r   r   r   r   r*   rU   joinr%   mapr  r   rD   r   r   rO   r   get_arg_specializationexec)sigkparamsbackendspecializationr-   kpr   r   r   r@   ry   	func_bodyr   func_namespaces                 r:   create_function_from_signaturer1    s    s~~#g,...N++-w7 0b??!!N4&":;!#v'H$&$8$8fJ!@@GfE$TF"XJbBugQOC!!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  %0N=!(>w?]?](^N$% 	N# .))% R
s    I=A Jc                8    | j                    d| j                   S )N.)rB   r   fns    r:   get_full_namer6    s    mm_Aboo.//r;   c                  ,    e Zd ZU ded<   ded<   ded<   y)JitFunctionInfor   rJ   rU   r-   rO   jit_functionN)rR   rB   r   r  r   r;   r:   r8  r8    s    
Ir;   r8  c                       e Zd Zd Z	 	 ddZd Zd Zd Zd Z	 	 ddZ	d Z
ed	        Zed
        Zd Zd Zd Zd Z fdZ fdZd Z xZS )rO   c                     y)NFr   r?   s    r:   is_gluonzJITFunction.is_gluon  s    r;   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   reprr5  compileis_manual_warmupalready_compiled)r5  r   rB   r&  r%  paramsr-   r@  rA  rB  rC  rD  r6  r  rE  r8  )r7   hookr   r
  r?  r  r  rF  rH  r-   rJ   r   r   	arg_reprsrI  	full_namerG  r   s                     r:   
_call_hookzJITFunction._call_hook  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v #" **((!,, ' 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)callablepre_run_hooksr   )r7   rN  s     r:   add_pre_run_hookzJITFunction.add_pre_run_hook  s$    
 ~~!!$'r;   c                    ddl m}m}m}m} t
        j                  j                         } ||      }|| _        || _        || _        t        | j                  | j                  |      }i |||fS )z1
        Precompute as much as possible.
        r   )CompiledKernelrJ  	ASTSourcemake_backend)compilerrW  rJ  rX  rY  r   activeget_current_targetr1  r
  rM  )r7   rW  rJ  rX  rY  r   r,  binders           r:   create_binderzJITFunction.create_binder  sb     	PO113v&,"/WU67F**r;   c                  |j                  d| j                        xs t        j                  j                  |d<   t        j
                  j                         }t        j
                  j                  |      }| j                  D ]
  } ||i |  | j                  |   \  }}	}
} ||i |\  }}}t        |      t        |      z   }|j                  |d       }||
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   	 }}t#        |d       }|D ci c]  }||
j+                  t%        ||              }}| j-                  t        j                  j.                  ||||||g|      ry | j1                  | |||      }| j3                  ||	|j                        }|||<   | j-                  t        j                  j4                  ||||||g|       t7               }| j8                  j;                         D ]6  \  \  }}\  }} | j                  ||      x}!|k7  s$t=        d| d| d|!        |s|J t?        |      r ||      }tA        |      }"|d   }#|"dkD  r|d   nd}$|"dkD  r|d   nd}% |jB                  ||g|j)                          }& |jD                  |#|$|%||jF                  |jH                  |&t        j                  jJ                  t        j                  jL                  g	|j)                           |S c c}w c c}w c c}}w c c}w c c}w c c}w )Ndebugr   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   )rX   r`   s     r:   r   z!JITFunction.run.<locals>.<lambda>G  s    sk?Q r;   r   c                "    t        |t              S r=   )r*   rU   )rX   r   s     r:   r   z!JITFunction.run.<locals>.<lambda>K  s    As9K r;   )r   r  rM   z1 has changed since we compiled this kernel, from z to r   )'r_   r`  r   runtimer   r[  get_current_deviceget_current_streamrT  device_cachesrU   parse_optionsrM  r-   r%  r  KeyErrorr   r   r%   r   
parse_attrrQ  jit_cache_hookrX  rJ  jit_post_compile_hookobjectr5   r  rQ   rS  r"   launch_metadatar  functionpacked_metadatalaunch_enter_hooklaunch_exit_hook)'r7   r  r  rx   r   r?  rb  rN  kernel_cacher   r,  r]  
bound_argsr-  r  r   kernelr   sigkeyssigvalsrV   vr
  
constexprspathattrvalsr  r8   not_presentr-   rX   r`   globals_dictnewVal	grid_sizegrid_0grid_1grid_2ro  s'                                          r:   r  zJITFunction.run#  s5    **Wdjj9PU]]=P=Pw 11311&9 && 	"D$!&!	" 150B0B60J-fgv /5d.Ef.E+
NG .!CL0!!#t, >++F3G'+{{3!qvv3G3%34qt4G4,/,AB&1aABIB .o0oo.6)e+ee)6)e+ee) ]G,,,'1A"#WZ[#[\\] 'w0QRJ_ijW[$ 1$z7H7H7J2KT RRjJj&45!5H5!(,KLETYZqQ**+<Xq+IJJZEZu}};;S)VU_ahkpjq%'..y*eDC\\#fg>N>N\OF &LOOEMM??iQWYcelotnu"$ 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_ 44B k5Zs$   =OO 3O%+O+O0-#O5c                T    | j                   | j                  S | j                  |      S r=   )_repr_fn_name)r7   rX   s     r:   rI  zJITFunction.reprm  s"     $

 2t}}E

1Er;   c	           	        |r|ng }|r|ng }|| _         |j                  | _        || _        t	        j
                  |      | _        || _        || _        t	        j                  |      d   | _	        || _
        t        |      | _        || _        g | _        t        | j
                  j                   j#                               D ]T  \  }	}
|	|v xs |
j$                  |v }|	|v xs |
j$                  |v }| j                  j'                  t)        |	|
||             V t+        j,                  t	        j.                  |            }|t1        j2                  d|t0        j4                        j7                         d  }| j9                  |       t;        | j<                        | _        d | _         i | _!        d | _"        || _#        || _$        | j                  D cg c]  }|j$                   c}| _%        | j                  D cg c]  }|jL                  s|jN                   c}| _(        g | _)        |jT                  | _*        |jV                  | _+        |jX                  | _,        |jZ                  | _-        |j                  | _        y c c}w c c}w )Nr   z^def\s+\w+\s*\().r5  rB   rJ   versionrD   r
  r   r   getsourcelinesstarting_line_numberr  r6  r  ro  rM  	enumerater$  r   r-   r   r   textwrapdedent	getsourceresearch	MULTILINEstart_unsafe_update_srcr   r^  rh  hashr5   rv  r`  rN   	arg_namesr   r   rz  rT  r   rR   r   __globals__)r7   r5  r  r   r   r`  rN   rI  ro  ir   dnsdns_oar8   ps                  r:   r,   zJITFunction.__init__p  s#   1B-Ki)Goq&mm **2.!2.L+$+$:$:2$>q$A!
%b).!$..";";"B"B"DE 	CHAu((KEJJ:K,KC88hEJJJh<hFKK{1eS&AB	C oog//34")).R\\BHHJKL$(););<	 TV 
  +/++6Q!&&6*.++HQ155H   zzOO>>-- 7Hs   J=JJc                n    | j                   t        j                  | j                        j                  z  S r=   )r  rD   getclosurevarsr5  r3   r?   s    r:   get_capture_scopezJITFunction.get_capture_scope  s(    '"8"8"A"K"KKKr;   c                   | j                   t        j                  | j                        j                  }t        | j                  | j                  || j                        }|j                  | j                                |j                  t        | j                        z   | _         t        t        |j                   j#                                     | _        | j                   S )N)r-   r2   r3   r8   )r  rD   r  r5  r3   r   r  r  r8   rn   parser@   rU   r  dictsortedr5   r  )r7   r3   dependencies_finders      r:   rT   zJITFunction.cache_key  s     99..tww7AAI"4$--QUQaQamv9=#C%%djjl3+//#d6O6O2PPDI$(0C0T0T0Z0Z0\)]$^D!yyr;   c                    ddl m} |S )Nr   r   )r   r   )r7   r   s     r:   rc   zJITFunction.type  s    2r;   c               \     | j                   t        t        j                  |      |dd|S )NTr  )r  r'  
MockTensor
wrap_dtype)r7   r  rx   r   s       r:   r  zJITFunction.warmup  s*    txxZ5J5JD1QT$\U[\\r;   c           	     8   ddl m}m} dd l}dd lm} t        j                  j                         }|j                  |      }|d   | j                  k7  rt        d|d    d| j                         t        t        |d         }|d   }	t        ||	      D 
ci c]4  \  }
}|
|j                  j!                  |      r|j                  |      n|6 }}
}t        t        |d	         }|d
   }t#        t        ||            }t#        |d   j%                               } || |||      }|d   j%                         D 
ci c]#  \  }
}|
t'        |t(              rt        |      n|% }}
}|d   }
 ||d |      }|| j*                  |   d   |
<   |S c c}}
w c c}}
w )Nr   )rJ  rX  r   r-   zSpecialization data is for z but trying to preload for r  r  r  r  r
  r  r   )rZ  rJ  rX  r  triton.languager   r   r[  rf  loadsr  rQ   r'  r   r%  r   is_dtyper  r  r*   r%   rh  )r7   rG  rJ  rX  r  tlr?  deserialized_objr  r  r   rr   r  r  r  r  r
  r8   r  rv  s                       r:   preloadzJITFunction.preload  s   1$113::&9:F#t}}4-.>v.F-GGbcgcpcpbqrt tE#3O#DE(9 "-?
U BHH$5$5e$<%%G
	 
  0 >?
%l3
SZ01)+6<<>?	iE: /y9??A
U E4!8ueC
 
 u%dG,-36"1%c*!

s   9F?(Fc                   t        j                  | j                        }t        |t         j                        sJ t        |j                        dk(  sJ t        |j                  d   t         j                        sJ |S )Nr   r   )re   r  r8   r*   Moduler"   bodyFunctionDef)r7   trees     r:   r  zJITFunction.parse  s_    yy"$

+++499~"""$))A,888r;   c                    t        d      )Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rQ   )r7   rx   r   s      r:   __call__zJITFunction.__call__  s    WXXr;   c                V    |dk(  rt        d| d      t        t        |   ||       y )Nr8   zCannot set attribute 'zX' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)AttributeErrorr+   rO   __setattr__)r7   r-   rr   r9   s      r:   r  zJITFunction.__setattr__  s<    5= #9$ @, "- . . 	k4,T59r;   c                4    d| _         t        | 	  d|       y)z
        The only method allowed to modify src.
        Bypasses the __setattr__ restriction by calling super().__setattr__ directly.
        Nr8   )r  r+   r  )r7   new_srcr9   s     r:   r  zJITFunction._unsafe_update_src  s    
 	E7+r;   c                P    d| j                    d| j                  j                   dS )NzJITFunction(:r  )rJ   r5  r   r?   s    r:   __repr__zJITFunction.__repr__  s&    dkk]!DGG,@,@+ACCr;   )r   zbool | None)NNNNNNN)rR   rB   r   r<  rQ  rU  r^  r  rI  r,   r  r   rT   rc   r  r  r  r  r  r  r  r   r   s   @r:   rO   rO     s    ,
 
,
\(+HTF mq;?<(|L 	 	  ]@Y:,Dr;   rO   c                     y r=   r   r4  s    r:   jitr    s    r;   r  rI  ro  r   r   r`  rN   c                     y r=   r   r  s          r:   r  r    s     r;   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`  rN   rI  ro  )rS  r   re  	interpretinterpreterr  rO   )	r5  r  r`  r   r   ro  rN   rI  r  s	     r:   	decoratorzjit.<locals>.decorator8  sj    ||==""8&r7N_Fdlq08tUdf f "3/M! /	 	r;   r5  r   r   zJITFunction[T]r   )	r5  r  rI  ro  r   r   r`  rN   r  s	    ``````` r:   r  r    s&    : & 
~} r;   c                  F    e Zd ZdZed        Zd Zed        Zed        Zy)r  zr
    Can be used in place of real tensors when calling:
        kernel.warmup(MockTensor(torch.float32), ...)
    c                l    | j                   j                  dk(  r| j                  dk(  rt        |       S | S )Nr   torch)r9   rR   rB   r  )ry   s    r:   r  zMockTensor.wrap_dtype]  s.    ==!!W,71Jc?"
r;   c                    || _         y r=   )r   )r7   r   s     r:   r,   zMockTensor.__init__c  s	    
r;   c                      yNr   r   r   r;   r:   r   zMockTensor.data_ptrf      r;   c                      yr  r   r   r;   r:   	ptr_rangezMockTensor.ptr_rangej  r  r;   N)	rR   rB   r   r   staticmethodr  r,   r   r  r   r;   r:   r  r  W  sH    
  
    r;   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 r=   )r   r   datar?  shape)r7   r   r   s      r:   r,   zTensorWrapper.__init__q  s5    
	II	kkYY__
r;   c                6    | j                   j                         S r=   )r   r   r?   s    r:   r   zTensorWrapper.data_ptrx  s    yy!!##r;   c                4     | j                   j                  | S r=   )r   stride)r7   rx   s     r:   r  zTensorWrapper.stride{  s    tyy&&r;   c                <    d| j                    d| j                   dS )NzTensorWrapper[r>  r  )r   r   r?   s    r:   __str__zTensorWrapper.__str__~  s    

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

#r;   c                ^    t        | j                  j                         | j                        S r=   )r  r   cloner   r?   s    r:   r  zTensorWrapper.clone  s    TYY__.

;;r;   c                `    t        | j                  j                  |      | j                        S r=   )r  r   tor   )r7   r?  s     r:   r  zTensorWrapper.to  s     TYY\\&14::>>r;   c                `    t        | j                  j                  |      | j                        S r=   )r  r   	new_emptyr   )r7   sizess     r:   r  zTensorWrapper.new_empty  s"    TYY007DDr;   Nr   )rR   rB   r   r,   r   r  r  r  r  r  r  r  r  r   r;   r:   r  r  o  s5    %$';(:$<?Er;   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 r3  )r*   r  r   r   r   r   rc   )r   r   s     r:   reinterpretr    sk    &-(FKK%%%;; !e44		$VU++/V~Q?@@r;   c                d   | }t        |t              s|j                  }t        |t              s|j                  j                  j                  }t        j                  |j                        \  }}t        |      D ].  \  }}|j                         j                  d      s&||z  } ||fS  ||fS )Nzdef )
r*   rO   r5  __code__co_filenamerD   r  r  r   rG   )r5  base_fn	file_namelines
begin_lineidxlines          r:   get_jit_fn_file_liner    s    G+.** +.

##//I..wzz:E: u% 	T::<""6*#Jj  	 j  r;   r   )Fr  )rI  Optional[Callable]ro  r  r   Optional[Iterable[int | str]]r   r  r`  Optional[bool]rN   r  r   zCallable[[T], JITFunction[T]]r=   )r5  zOptional[T]rI  r  ro  r  r   r  r   r  r`  r  rN   r  r   z4Union[JITFunction[T], Callable[[T], JITFunction[T]]])?
__future__r   r   re   ri   r.   rD   r   r  r  collectionsr   dataclassesr   	functoolsr   typingr   r	   r
   r   r   r   r   r   r   r   triton.tools.tensor_descriptorr   typesr   rC   r   runtime.driverr   _utilsr   r   r   r   rR   r"   rH   r   NodeVisitorr   r   r   r   r   r   r   r   r  r1  r6  r8  rO   r  r  r  r  r  r   r;   r:   <module>r     s   , 
     	  # ! % d d d ;   # e e.3~../CLH! H!`D4/> />d 	 6r@	Zgaj 	Z	7*t0   eD/!$ eDZ	 
 
 
 #*.7;DH #
 
 (	

 5
 %B
 
 
 #
 

 4 #*.7;DH #44 	4
 (4 54 %B4 4 4 :4x 0"E "EJA!r;   