o
    ۗi                  	   @  sL  d dl mZmZ d dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
mZ d dlmZ d dlmZmZmZmZmZmZmZmZmZmZ ddlmZ d dlmZ eded	  Zed
Z G dd dej!Z"d_ddZ#G dd dZ$dd Z%i Z&d`ddZ'G dd dee  Z(dd Z)dd Z*i ddd d!d"d#d$d%d&d!d'd(d)d(d*d#d+d,d-d,d.d/d0d1d2d3d4d5d6d7d8d9d:d;d<d=d>d?d@dAZ+e,e+- D ]Z.e.e+e.< qG dBdC dCe(e  Z/edadFdGZ0eddddddddHdbdSdGZ0	dcddddddddHdddVdGZ0G dWdX dXZ1G dYdZ dZZ2d[d\ Z3d]d^ Z4dS )e    )annotationsdivisionN)defaultdict)cached_property)
CallableGenericIterableOptionalTypeVarUnionoverloadDictAnyTuple   )driver)
ModuleTypez.runtime.jitTc                      s   e Zd ZdZd fddZedd Zdd	 Zd
d Zdd Z	dd Z
dd Zdd Zdd Zdd Zdd Zdd Zdd Z  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                   sB   t    || _t|d| _|| _h d| _i | _	d| _
d S )Nutf-8>
   intlenmaxminlistfloatprintrangegetattr
isinstanceF)super__init__namehashlibsha256encodehasherglobalssupported_python_builtinsused_global_valsvisiting_arg_default_value)selfr$   r)   src	__class__ P/home/ubuntu/transcripts/venv/lib/python3.10/site-packages/triton/runtime/jit.pyr#   $   s   


zDependenciesFinder.__init__c                 C  
   | j  S N)r(   	hexdigestr-   r1   r1   r2   retH      
zDependenciesFinder.retc                 C  s&   t |jrdS t|dd}|tS )NT
__module__ )inspect	isbuiltinfuncr    
startswithTRITON_MODULE)r-   noder=   moduler1   r1   r2   _is_triton_builtinL   s   
z%DependenciesFinder._is_triton_builtinc                 C  s   t |tr]| j |j @ D ].}|\}}| j| \}}|j| \}}||kr=td| d| d| j d|j d| dq| j|j |j}|t	t
|dd7 }| j|d	 d S d S )
NGlobal variable z has value z when compiling z, but inner kernel z has conflicting value z7 from when it was first compiled.  This is not allowed.noinlineFr   )r!   JITFunctionr+   keysRuntimeErrorr$   __name__update	cache_keystrr    r(   r'   )r-   r=   kvar_name_v1v2func_keyr1   r1   r2   _update_hashR   s   
&zDependenciesFinder._update_hashc                 C  s   t |jtju r|jS |j| jv rd S | j|jd }|d urG| jsGt |t	urGt
|tsGt|ddsG|j| jvrG|| jf| j|jt| jf< | | |S )N__triton_builtin__F)typectxastStoreidlocal_namesr)   getr,   r   r!   rE   r    r*   r+   rR   )r-   r@   valr1   r1   r2   
visit_Named   s"   

zDependenciesFinder.visit_Namec                   s    fdd|j D S )Nc                   s   g | ]}  |qS r1   )visit).0eltr6   r1   r2   
<listcomp>       z2DependenciesFinder.visit_Tuple.<locals>.<listcomp>)eltsr-   r@   r1   r6   r2   visit_Tuple   s   zDependenciesFinder.visit_Tuplec                 C  sf   |  |j}t|tjr|  |j}t|tjs|d u s$t|ddtkr&d S t||j}| | |S )NrH   r:   )	r]   valuer!   rV   	Attributer    r?   attrrR   )r-   r@   lhsr7   r1   r1   r2   visit_Attribute   s   
z"DependenciesFinder.visit_Attributec                 C  s"   dd |j j D | _| | d S )Nc                 S  s   h | ]}|j qS r1   arg)r^   rk   r1   r1   r2   	<setcomp>       z7DependenciesFinder.visit_FunctionDef.<locals>.<setcomp>)argsrY   generic_visitrc   r1   r1   r2   visit_FunctionDef   s   z$DependenciesFinder.visit_FunctionDefc                   sn    fdd}t |j|j|jr|jgng |jD ]} | q||j |jd ur0 |j ||j	 d S )Nc                   sB   z j rJ d _ | D ]}|d ur | qW d _ d S d _ w )NTF)r,   r]   )defaultsexprr6   r1   r2   visit_defaults   s   

z:DependenciesFinder.visit_arguments.<locals>.visit_defaults)
	itertoolschainposonlyargsrn   vararg
kwonlyargsr]   kw_defaultskwargrq   )r-   r@   rs   rk   r1   r6   r2   visit_arguments   s   (


z"DependenciesFinder.visit_argumentsc                 C  s:   |  |}t|tr|  jt|O  _d S | j| d S r4   )r]   r!   r   rY   setadd)r-   r@   targetr1   r1   r2   visitAssnTarget   s   

z"DependenciesFinder.visitAssnTargetc                 C  s4   t |jdkrtd| |jd  | | d S )N   z2Simultaneous multiple assignment is not supported.r   )r   targets	TypeErrorr   ro   rc   r1   r1   r2   visit_Assign   s   zDependenciesFinder.visit_Assignc                 C     |  |j | | d S r4   r   r~   ro   rc   r1   r1   r2   visit_AnnAssign      z"DependenciesFinder.visit_AnnAssignc                 C  r   r4   r   rc   r1   r1   r2   	visit_For   r   zDependenciesFinder.visit_For)r   r   )rH   r9   __qualname____doc__r#   propertyr7   rB   rR   r\   rd   ri   rp   r{   r   r   r   r   __classcell__r1   r1   r/   r2   r      s     $

 	r   r   rK   c                 C  s&   t | tr| jS t | tr| S t| S r4   )r!   rT   rH   rK   repr)tyr1   r1   r2   _normalize_ty   s
   

r   c                   @  sn   e Zd ZdZdd	d
Zedd Zedd Zedd Zedd Z	edd Z
edd Zed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                 C  s   || _ || _|| _|| _d S r4   )r   _paramr   r   )r-   r   r   r   r   r1   r1   r2   r#      s   
zKernelParam.__init__c                 C     | j jS r4   )r   r$   r6   r1   r1   r2   r$         zKernelParam.namec                 C  s(   | j jr| j jtjjkrdS t| j jS )Nr:   )r   
annotationr;   	Parameteremptyr   r6   r1   r1   r2   r      s   zKernelParam.annotationc                 C  sZ   | j }dD ]\}}|||t| d  }|r$||v r$| |   S q|dkr+dS dS )N))uintu)r   ir   u1r:   )r   findr   )r-   r   ty1ty2widthr1   r1   r2   annotation_type   s   zKernelParam.annotation_typec                 C  s
   d| j v S )N	constexpr)r   r6   r1   r1   r2   is_constexpr  r8   zKernelParam.is_constexprc                 C  s   d| j v o| j S )Nconst)r   r   r6   r1   r1   r2   is_const	  s   zKernelParam.is_constc                 C  r   r4   )r   defaultr6   r1   r1   r2   r     r   zKernelParam.defaultc                 C  s   | j jtjjkS r4   )r   r   r;   r   r   r6   r1   r1   r2   has_default  s   zKernelParam.has_defaultN)r   r   r   r   r   r   r   r   )rH   r9   r   r   r#   r   r$   r   r   r   r   r   r   r   r1   r1   r1   r2   r      s"    







r   c                 C  sP   |rt | dr|  d dkrdS t| tr&|r | d dkr dS | dkr&dS dS )Ndata_ptr   r   Dr   1N)hasattrr   r!   r   )valignr1   r1   r2   compute_spec_key  s   
r   Fc                 C  s   | d u rdS t | trdS t | tr(d| kr| dkrdS d| kr&| dkr&dS d	S t | tr/d
S t| dr6dS | j|f}t|d }|d u r^|d rKdndtt	|d 
dd   }|t|< |S )Nnonei1   i32                u64i64fp32tma_desc_cpu_ptr	nvTmaDescr   *k*r   .)r!   r   r   r   r   dtype	dtype2strrZ   type_canonicalisation_dictrK   split)rk   r   dskresr1   r1   r2   mangle_type&  s(   




*r   c                   @  s    e Zd ZU ded< dddZdS )KernelInterfacer   runr   c                   s    fd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                    s   j |  dd|S )NFgridwarmup)r   )rn   kwargsr   r-   r1   r2   <lambda>J  ra   z-KernelInterface.__getitem__.<locals>.<lambda>r1   )r-   r   r1   r   r2   __getitem__D  s   zKernelInterface.__getitem__N)r   r   )rH   r9   r   __annotations__r   r1   r1   r1   r2   r   A  s   
 r   c           	      C  s@   dd |  D }dd l}| ||| |j|d}||}|S )Nc                 S  s*   i | ]\}}||j jd krt|n|qS r   )r0   rH   rK   r^   keyre   r1   r1   r2   
<dictcomp>O  s   * z1serialize_specialization_data.<locals>.<dictcomp>r   )r$   	signature	constantsattrsoptionsr   )itemsjsonto_dict__dict__dumps)	r$   r   r   r   r   r   r   objserialized_objr1   r1   r2   serialize_specialization_dataN  s   
r   c                 C  s  t | jt |ksJ g }g }g }g }g }g }t| j |D ]p\\}	}
}|
jtjju r=||	 |d|	 d|	  n||	 d|	  |d|	 d|	  |j	r[||	 q||	 |j
su|jsn|d|	  n|d|	  |jr|d|j  q|d|	|jrdnd	f  qd
dd || D }d
dd |D }d
dd |D }|d d|}d|}d|||||f }dd | j D }t|d< |j|d< t|| |d S )a2  
    Equivalent to sig.bind followed by apply_defaults. This generates a
    native Python function (using exec) which can be memoized on a per-kernel
    basis to avoid having to run these expensive functions -- which constitute
    much of the kernel launch overhead -- every time we run the kernel.
    'z': z	=default_z compute_spec_key(%s, align=True)z!compute_spec_key(%s, align=False)z"%s"zmangle_type(%s, %s)TrueFalser:   c                 S     g | ]}|d  qS , r1   r^   xr1   r1   r2   r`         z2create_function_from_signature.<locals>.<listcomp>c                 S  r   r   r1   r   r1   r1   r2   r`     r   c                 S  r   r   r1   r   r1   r1   r2   r`     r   z**excess_kwargsr   zFdef dynamic_func(%s):
    return {%s}, (%s), (%s), (%s), excess_kwargsc                 S  s,   i | ]\}}|j tjjurd | |j qS )default_)r   r;   r   r   )r^   r$   r   r1   r1   r2   r     s
    z2create_function_from_signature.<locals>.<dictcomp>r   r   dynamic_func)r   
parameterszipr   r   r;   r   r   appendr   r   r   r   r   joinr   r   exec)sigkparamsbackend	func_argsdict_entriesconstexpr_valsnon_constexpr_valssignature_typesspecialisationsr$   spkprJ   args_strdict_str	func_bodyfunc_namespacer1   r1   r2   create_function_from_signatureY  sN   







r  r   r   
float8e4nvfp8e4nvfloat8e5fp8e5float8e4b15fp8e4b15float8_e4m3fn
float8e4b8fp8e4b8float8_e4m3fnuzfloat8_e5m2float8e5b16fp8e5b16float8_e5m2fnuzfloat16fp16bfloat16bf16float32r   float64fp64int8i8int16i16int32r   r   u8u16u32r   )int64uint8uint16uint32uint64c                      s   e Zd ZdZdZedd Zed!ddZdd Zd	d
 Z	dd Z
dd Zdd Z		d"ddZedd Zdd Zdd Zdd Zdd Z fddZdd  Z  ZS )#rE   Nc                 C  s   t | dr| jS t| trdS t| tr*d| kr| dkrdS d| kr(| dkr(dS d	S t| tr1d
S | d u r7d S tdt|  d|  )Nr   r   r   r   r   r   r   r   r   r   zUnsupported type z for )r   r   r!   r   r   r   r   rT   rj   r1   r1   r2   _key_of  s   



zJITFunction._key_ofFc                 C  sH   | d u rdS t | tr| S t| dd }t| }|rdnd}|| S )N*i8r   r   r   r   )r!   rK   r   r   )r   r   	dtype_str	const_strr1   r1   r2   _type_of  s   
zJITFunction._type_ofc                 C  s   t t| j|}|S r4   )dictr   
constexprs)r-   constexpr_keyr   r1   r1   r2   _make_constants  s   zJITFunction._make_constantsc	                 C  s   |rt jnt j}	|	d u rdS | jj}
| jj}ddd t| j|d D }|
 d|j	 d|j
 d|j d	|j d
| d}G dd d}t|
|||d ||}||||j	|j
|j|j|j|||d}|	|||||
| d|i||ddS )NFr   c                 S  s    g | ]\}}|j  d | qS )z: r$   )r^   r   r   r1   r1   r2   r`     s     z*JITFunction._call_hook.<locals>.<listcomp>r   z[num_warps=z, num_ctas=z, num_stages=z, enable_fp_fusion=]()c                   @  s   e Zd Zdd ZdS )z/JITFunction._call_hook.<locals>.JitFunctionInfoc                 S  s   || _ || _|| _d S r4   )rA   r$   jit_function)r-   rA   r$   r4  r1   r1   r2   r#     s   z8JITFunction._call_hook.<locals>.JitFunctionInfo.__init__N)rH   r9   r   r#   r1   r1   r1   r2   JitFunctionInfo  s    r5  r   )r   devicer   	num_warpsnum_ctas
num_stagesenable_fp_fusionextern_libsconfigsspecialization_data	is_warmupr   )r   r   fncompileis_manual_warmupalready_compiled)rE   
cache_hookcompiled_hookr?  rH   r9   r   r   paramsr7  r8  r9  r:  r   r;  )r-   r   r   r6  r   r   r<  r>  beforehookr$   rA   	arg_reprsr   r5  r=  r   r1   r1   r2   
_call_hook  s:    0

zJITFunction._call_hookc                 C  s   t |sJ | j| dS )z
        Add a hook that will be executed prior to the execution of run
        function with args and kwargs passed into the kernel
        N)callablepre_run_hooksr   )r-   rG  r1   r1   r2   add_pre_run_hook  s   zJITFunction.add_pre_run_hookc                 C  s   ddl m}m}m}m} || _|| _|| _|| _t| j| j|| _dd t	| jD | _
dd t	| jD | _dd t	| jD | _dS )z1
        Precompute as much as possible.
        r   )CompiledKernelr@  	ASTSourcemake_backendc                 S  s   g | ]	\}}|j r|qS r1   r   r^   r   pr1   r1   r2   r`   -      z-JITFunction.create_binder.<locals>.<listcomp>c                 S  s   g | ]	\}}|j s|qS r1   rP  rQ  r1   r1   r2   r`   .  rS  c                 S  s    g | ]\}}|j s|js|qS r1   )r   r   rQ  r1   r1   r2   r`   /  s    
N)compilerrM  r@  rN  rO  r  r   rE  binder	enumerateconstexpr_indicesnon_constexpr_indicesspecialised_indices)r-   r   rM  r@  rN  rO  r1   r1   r2   create_binder#  s   zJITFunction.create_binderc          (   
     s6  | ddptj dddk|d< ddlm} tj }tj|}tj	 }||}	j
D ]	}
|
|i | q.jd u rB|	 j|i |\}}}}}d|t||f }j|  |d }|d u r|	|}d	|vsuJ d
d|vs}J dd|vsJ d|D ]}||jvrtd| qt| }fddjD }|d t| }dd t||D }|	j|f}|d    fddt|jD }| D ]\}}t|rtd| dqֈj|||||||ddrd S |||d }j |||jd}|j| |< j|||||||dd t! }j" D ]!\\}}\} }!|! || }"| krEt#d| d|  d|" q%|s|d usQJ t|rZ||}t|}#|d }$|#dkrk|d nd}%|#dkrv|d nd}&|j$||g|R  }'|j%|$|%|&||j&|j'|'j(j)j(j*g	|R   |S )NdebugFTRITON_DEBUG0r   r   )rO  r:   device_typez=device_type option is deprecated; current target will be usedr6  z8device option is deprecated; current device will be usedstreamz8stream option is deprecated; current stream will be usedz2Keyword argument %s was specified but unrecognisedc                   s   g | ]} j | jqS r1   )rE  r$   )r^   r   r6   r1   r2   r`   \      z#JITFunction.run.<locals>.<listcomp>c                 S  s"   i | ]\}}||d krdn|qS )r   r)  r1   )r^   rL   r   r1   r1   r2   r   ^  s   " z#JITFunction.run.<locals>.<dictcomp>r   c                   s0   i | ]\}}|j s|j v s|d u r|j|qS r4   )r   r   r$   )r^   r   rR  )constant_paramsr1   r2   r   b  s    zCallable constexpr at index z is not supportedT)rF  )r~   r   rC   z1 has changed since we compiled this kernel, from z to r   )+rZ   osenvironrT  rO  r   activeget_current_deviceget_current_streamget_current_targetrK  rU  rZ  r   rK   cacheparse_optionsr   KeyErrortuplevaluesrX  r   r   get_attrs_descriptorrE  get_constantsr   rJ  r   rI  rN  r@  objectr+   rG   launch_metadatar   functionpacked_metadatarM  launch_enter_hooklaunch_exit_hook)(r-   r   r   rn   r   rO  r6  r_  r~   r   rG  
bound_argssig_and_specr   r   excess_kwargsr   kernelr   rL   
bound_valssigkeyssigvalsr   r<  r   r   rk   r.   not_presentr$   rN   r[   globals_dictnewVal	grid_sizegrid_0grid_1grid_2rp  r1   )ra  r-   r2   r   3  s   "










zJITFunction.runc	                   sn  |r|ng }|r
|ng } | _  j| _|| _t | _|| _|| _t d | _	 fdd| _
|| _d | _g | _t| jj D ]!\}	}
|	|v pO|
j|v }|	|v pX|
j|v }| jt|	|
|| qCtt | _| jtd| jtj d  | _tt| _d | _i | _ d | _!|| _"dd | jD | _#dd | jD | _$g | _% j&| _& j'| _' j(| _( j| _d S )Nr   c                   s   d u r j S | S r4   )rH   )rN   r?  r   r1   r2   r     ra   z&JITFunction.__init__.<locals>.<lambda>z^def\s+\w+\s*\(c                 S  s   g | ]}|j qS r1   r1  r^   rR  r1   r1   r2   r`     rm   z(JITFunction.__init__.<locals>.<listcomp>c                 S  s   g | ]}|j r|jqS r1   )r   r   r  r1   r1   r2   r`     r`  ))r?  r9   rA   versionr;   r   r   r   getsourcelinesstarting_line_numberr   rp  rU  rE  rV  r   rl  r$   r   r   textwrapdedent	getsourcer.   research	MULTILINEstartr   r-  rh  hashr+   rx  rD   	arg_namesr.  rK  r   rH   __globals__)r-   r?  r  r   r   r[  rD   r   rp  r   r   dnsdns_oar1   r  r2   r#     s>   "
zJITFunction.__init__c                 C  sX   | j d u r)t| j| j| jd}||   |jt| j	 | _ t
t|j | _| j S )N)r$   r)   r.   )r  r   rH   r  r.   r]   parser7   rK   r  r-  sortedr+   r   )r-   dependencies_finderr1   r1   r2   rJ     s   
zJITFunction.cache_keyc                O  s   | j ttj||dd|S )NTr   )r   map
MockTensor
wrap_dtype)r-   r   rn   r   r1   r1   r2   r     s   zJITFunction.warmupc                   s   ddl m}m} ddlm} dd l}dd lm  tj	
 }||}|d | jjkr8td|d  d| jj  fdd	|d
  D }t|d  }	|| |	|||d }
dd	 |d  D }|d }||
d |}|| j| |< |S )Nr   )r@  rN  r   )AttrsDescriptorr$   zSpecialization data is for z but trying to preload for c                   s,   i | ]\}}| j |r  |n|qS r1   )r   is_dtyper   tlr1   r2   r     s    z'JITFunction.preload.<locals>.<dictcomp>r   r   r   c                 S  s(   i | ]\}}|t |trt|n|qS r1   )r!   r   rk  r   r1   r1   r2   r     s    r   r   )rT  r@  rN  triton.backends.compilerr  r   triton.languagelanguager   rd  re  loadsr?  rH   rG   r   r-  	from_dictrh  )r-   r=  r@  rN  r  r   r6  deserialized_objr   r   r.   r   r   rx  r1   r  r2   preload  s,   




zJITFunction.preloadc                 C  sH   t | j}t|t jsJ t|jdksJ t|jd t js"J |S )Nr   r   )rV   r  r.   r!   Moduler   bodyFunctionDef)r-   treer1   r1   r2   r    s
   zJITFunction.parsec                 O  s   t d)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rG   )r-   rn   r   r1   r1   r2   __call__  s   zJITFunction.__call__c                   s(   t t| || |dkrd | _d S d S )Nr.   )r"   rE   __setattr__r  )r-   r$   re   r/   r1   r2   r     s   
zJITFunction.__setattr__c                 C  s   d| j  d| jj dS )NzJITFunction(:r3  )rA   r?  rH   r6   r1   r1   r2   __repr__  s   zJITFunction.__repr__F)NNNNNNN)rH   r9   r   rC  rD  staticmethodr(  r,  r0  rI  rL  rZ  r   r#   r   rJ   r   r  r  r  r  r  r   r1   r1   r/   r2   rE     s.    
5^
<
	rE   r?  JITFunction[T]c                 C     d S r4   r1   )r?  r1   r1   r2   jit     r  r  r   rp  r   r   r[  rD   r   Optional[Callable]rp  r   Optional[Iterable[int]]r   r[  Optional[bool]rD   Callable[[T], JITFunction[T]]c                 C  r  r4   r1   r  r1   r1   r2   r    s   Optional[T]4Union[JITFunction[T], Callable[[T], JITFunction[T]]]c          	        s.   d fdd}| dur|| 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
    r?  r   r   r  c              
     sX   t | sJ tdddkr ddlm} ||  dS t|  dS )NTRITON_INTERPRETr]  r   r   )InterpretedFunction)r  r   r   r[  rD   r   rp  )rJ  rb  getenvinterpreterr  rE   )r?  r  r[  r   r   rp  rD   r   r  r1   r2   	decorator@  s"   zjit.<locals>.decoratorNr?  r   r   r  r1   )	r?  r  r   rp  r   r   r[  rD   r  r1   r  r2   r  #  s   c                   @  s<   e Zd ZdZedd Zdd Zedd Zedd	 Zd
S )r  zr
    Can be used in place of real tensors when calling:
        kernel.warmup(MockTensor(torch.float32), ...)
    c                 C  s"   | j jdkr| jdkrt| S | S )Nr   torch)r0   rH   r9   r  rj   r1   r1   r2   r  e  s   zMockTensor.wrap_dtypec                 C  s
   || _ d S r4   r   )r-   r   r1   r1   r2   r#   k     
zMockTensor.__init__c                   C     dS Nr   r1   r1   r1   r1   r2   r   n  r  zMockTensor.data_ptrc                   C  r  r  r1   r1   r1   r1   r2   	ptr_ranger  r  zMockTensor.ptr_rangeN)	rH   r9   r   r   r  r  r#   r   r  r1   r1   r1   r2   r  _  s    

r  c                   @  sV   e Zd Zdd Zdd Zdd Zdd	d
Zdd Zdd Zdd Z	dd Z
dd ZdS )TensorWrapperc                 C  s*   || _ || _|j| _|j| _| jj| _d S r4   )r   basedatar6  shape)r-   r  r   r1   r1   r2   r#   y  s
   zTensorWrapper.__init__c                 C  r3   r4   )r  r   r6   r1   r1   r2   r     r  zTensorWrapper.data_ptrc                 C  s   | j |S r4   )r  stride)r-   r   r1   r1   r2   r    s   zTensorWrapper.strider   rK   c                 C  s   d| j  d| j dS )NzTensorWrapper[r2  r3  )r   r  r6   r1   r1   r2   __str__  s   zTensorWrapper.__str__c                 C  r3   r4   )r  element_sizer6   r1   r1   r2   r    r  zTensorWrapper.element_sizec                 C     t | j | jS r4   )r  r  cpur   r6   r1   r1   r2   r       zTensorWrapper.cpuc                 C  s   | j |j  d S r4   )r  copy_)r-   otherr1   r1   r2   r    r  zTensorWrapper.copy_c                 C  r  r4   )r  r  cloner   r6   r1   r1   r2   r    r  zTensorWrapper.clonec                 C  s   t | j|| jS r4   )r  r  tor   )r-   r6  r1   r1   r2   r    s   zTensorWrapper.toNr   rK   )rH   r9   r   r#   r   r  r  r  r  r  r  r  r1   r1   r1   r2   r  w  s    
r  c                 C  sP   t | tr|| jjkr| jS t| j|S t| drt| |S tdt|  d)Nr   zCannot reinterpret a r   )r!   r  r  r   r   r   rT   )tensorr   r1   r1   r2   reinterpret  s   


r  c                 C  sr   | }t |ts|j}t |tr|jjj}t|j\}}t|D ]\}}| 	dr4||7 } ||fS q ||fS )Nzdef )
r!   rE   r?  __code__co_filenamer;   r  rV  stripr>   )r?  base_fn	file_namelines
begin_lineidxliner1   r1   r2   get_jit_fn_file_line  s   


r  r  r  r  )r   r  rp  r  r   r  r   r  r[  r  rD   r  r   r  r4   )r?  r  r   r  rp  r  r   r  r   r  r[  r  rD   r  r   r  )5
__future__r   r   rV   r%   r;   rt   rb  r  r  collectionsr   	functoolsr   typingr   r   r   r	   r
   r   r   r   r   r   runtime.driverr   typesr   rH   r   r?   r   NodeVisitorr   r   r   r   r   r   r   r   r  r   r   rl  r   rE   r  r  r  r  r  r1   r1   r1   r2   <module>   s    0 
G0
D	

  X<"