o
    ohL                  	   @  s`  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 dd	lmZ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'i Z(g Z)dd Z*dad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/d0d1d2d3d4d5d6d7d8d9d:d;d<d=d>d?d@dAdBZ/e0e/1 D ]Z2e2e/e2< qG dCdD dDe,e# Z3edbdGdHZ4eddddddddIdcdTdHZ4	ddddddddddIdedWdHZ4G dXdY dYZ5G dZd[ d[Z6d\d] Z7d^d_ Z8dS )f    )annotationsdivisionN)defaultdict)cached_property)
CallableGenericIterableOptionalTypeVarUnionoverloadDictAnyTuple   )driver)
ModuleType)find_paths_ifget_iterable_pathz.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__ f/var/www/html/construction_image-detection-poc/venv/lib/python3.10/site-packages/triton/runtime/jit.pyr%   %   s   


zDependenciesFinder.__init__c                 C  
   | j  S N)r*   	hexdigestr/   r3   r3   r4   retI      
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?   moduler3   r3   r4   _is_triton_builtinM   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_keyr3   r3   r4   _update_hashS   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#   rG   r"   r,   r-   rT   )r/   rB   valr3   r3   r4   
visit_Namee   s"   

zDependenciesFinder.visit_Namec                   s    fdd|j D S )Nc                   s   g | ]}  |qS r3   )visit).0eltr8   r3   r4   
<listcomp>       z2DependenciesFinder.visit_Tuple.<locals>.<listcomp>)eltsr/   rB   r3   r8   r4   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 )NrJ   r<   )	r_   valuer#   rX   	Attributer"   rA   attrrT   )r/   rB   lhsr9   r3   r3   r4   visit_Attribute   s   
z"DependenciesFinder.visit_Attributec                 C  s"   dd |j j D | _| | d S )Nc                 S  s   h | ]}|j qS r3   arg)r`   rm   r3   r3   r4   	<setcomp>       z7DependenciesFinder.visit_FunctionDef.<locals>.<setcomp>)argsr[   generic_visitre   r3   r3   r4   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exprr8   r3   r4   visit_defaults   s   

z:DependenciesFinder.visit_arguments.<locals>.visit_defaults)
	itertoolschainposonlyargsrp   vararg
kwonlyargsr_   kw_defaultskwargrs   )r/   rB   ru   rm   r3   r8   r4   visit_arguments   s   (


z"DependenciesFinder.visit_argumentsc                 C  s:   |  |}t|tr|  jt|O  _d S | j| d S r6   )r_   r#   r   r[   setadd)r/   rB   targetr3   r3   r4   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   rq   re   r3   r3   r4   visit_Assign   s   zDependenciesFinder.visit_Assignc                 C     |  |j | | d S r6   r   r   rq   re   r3   r3   r4   visit_AnnAssign      z"DependenciesFinder.visit_AnnAssignc                 C  r   r6   r   re   r3   r3   r4   	visit_For   r   zDependenciesFinder.visit_For)r   r   )rJ   r;   __qualname____doc__r%   propertyr9   rD   rT   r^   rf   rk   rr   r}   r   r   r   r   __classcell__r3   r3   r1   r4   r      s     $

 	r   r   rM   c                 C  s&   t | tr| jS t | tr| S t| S r6   )r#   rV   rJ   rM   repr)tyr3   r3   r4   _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 r6   )r   _paramr   r   )r/   r   r   r   r   r3   r3   r4   r%      s   
zKernelParam.__init__c                 C     | j jS r6   )r   r&   r8   r3   r3   r4   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   r8   r3   r3   r4   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widthr3   r3   r4   annotation_type   s   zKernelParam.annotation_typec                 C  s
   d| j v S N	constexpr)r   r8   r3   r3   r4   is_constexpr  r:   zKernelParam.is_constexprc                 C  s   d| j v o| j S )Nconst)r   r   r8   r3   r3   r4   is_const
  s   zKernelParam.is_constc                 C  r   r6   )r   defaultr8   r3   r3   r4   r     r   zKernelParam.defaultc                 C  s   | j jtjjkS r6   )r   r   r=   r   r   r8   r3   r3   r4   has_default  s   zKernelParam.has_defaultN)r   r   r   r   r   r   r   r   )rJ   r;   r   r   r%   r   r&   r   r   r   r   r   r   r   r3   r3   r3   r4   r      s"    







r   c                     s6   t rt d S ddlm  d fdd	t  S )	Nr   )r   FTc                   s   d u rdS t  trd jfS t  rd fS t  tr dS t  trT|r. d|dnd } dkr8|r8dS d krD d	krDd
|fS d krP dkrPd|fS d|fS t  tr[dS t drbdS t  trfdd D } fdd}|dd |D }|dd |D }	||	fS  j|f}
t	
|
d }|d u r|
d rdndtt|
d dd   }|t	|
< |r d|dnd }||fS )N)r   Nr   )i1Nr   )alignr   )r   r   i   ii32l            l    u64i64)fp32Ntma_desc_cpu_ptr)	nvTmaDescNc                   s   g | ]}| qS r3   r3   r`   x)specialize_extraspecialize_implr3   r4   rb   :  rc   zCcreate_specialize_impl.<locals>.specialize_impl.<locals>.<listcomp>c                   s   t  drt |  S t| S )N_fields)hasattrrV   tuple)valsrl   r3   r4   <lambda>;      zAcreate_specialize_impl.<locals>.specialize_impl.<locals>.<lambda>c                 S     g | ]}|d  qS r   r3   r   r3   r3   r4   rb   <      c                 S  r   r   r3   r   r3   r3   r4   rb   =  r   z*k*r   .r   tensor)r#   rG   rL   r   r   r   r   r   dtype	dtype2strr\   type_canonicalisation_dictrM   split)rm   r   r   specialize_valuer   keyspec
make_tupletysrH   dskresr   r   )rm   r   r4   r   !  sD   








*z/create_specialize_impl.<locals>.specialize_impl)FTT)specialize_impl_cachelanguager   appendr3   r3   r   r4   create_specialize_impl  s   
(r   Fc                 C  s   t  }|| dd |dd S )Nc                 [     d S r6   r3   )rP   kwargsr3   r3   r4   r   O  s    zmangle_type.<locals>.<lambda>)r   r   )r   )rm   
specializer   r3   r3   r4   mangle_typeM  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   )rp   r   r   r/   r3   r4   r   [  rc   z-KernelInterface.__getitem__.<locals>.<lambda>r3   )r/   r   r3   r   r4   __getitem__U  s   zKernelInterface.__getitem__N)r   r   )rJ   r;   r   __annotations__r   r3   r3   r3   r4   r   R  s   
 r   c           	   	   C  sl   dd |  D }dd l}| |dd | D t| dd | D t| |j|d}||}|S )Nc                 S  s*   i | ]\}}||j jd krt|n|qS r   )r2   rJ   rM   r`   r   rg   r3   r3   r4   
<dictcomp>`  s   * z1serialize_specialization_data.<locals>.<dictcomp>r   c                 S     g | ]}t |qS r3   r   r   r3   r3   r4   rb   c  r   z1serialize_specialization_data.<locals>.<listcomp>c                 S  r   r3   r   r   r3   r3   r4   rb   d  r   )r&   	signatureconstant_keysconstant_vals
attrs_keys
attrs_valsoptionsr   )itemsjsonrH   r   values__dict__dumps)	r&   r   	constantsattrsr   r   r   objserialized_objr3   r3   r4   serialize_specialization_data_  s   $
r   c              
   C  s^  t | jt |ksJ g }t| j |D ]L\}}|jr&|d| d q|jr+dnd}|jr2dnd}|jr9dnd}d| d| d| d| d	}	|j	r[|d|j	 d	|	 d
 q||	  qdd }
dd
tt|
| j dg  dd
dd | j D  dd
| d}dd | j D }t|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("constexpr", )TrueFalsezspecialize_impl(z, specialize_extra, , z("z",) + z[1:]c                 S  s0   | d j tjju r| d S | d  d| d  S )Nr   r   z	=default_r   r=   r   r   )r   r3   r3   r4   r     s   0 z0create_function_from_signature.<locals>.<lambda>z
def dynamic_func(z	**optionsz):
    params = {c                 S  s   g | ]
}d | d| qS )'z': r3   )r`   r&   r3   r3   r4   rb     s    z2create_function_from_signature.<locals>.<listcomp>z}
    specialization = [,z-]
    return params, specialization, options
c                 S  s,   i | ]\}}|j tjjurd | |j qS )default_r   )r`   r&   r   r3   r3   r4   r     s
    z2create_function_from_signature.<locals>.<dictcomp>rG   r   r   dynamic_func)r   
parametersziprH   r   r   r   r   r   r   joinr   mapr   rG   r   get_arg_specializationexec)sigkparamsbackendspecializationr&   kpr   r   r   r9   rm   	func_bodyfunc_namespacer3   r3   r4   create_function_from_signaturek  s8   


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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 fddZdd Z  ZS )rG   Nc	                 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
|j d| d}G dd d}t|
|||d ||}||||j	|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   r3   r3   r4   rb          z*JITFunction._call_hook.<locals>.<listcomp>r   z[num_warps=z, num_ctas=z, num_stages=z, enable_fp_fusion=z, launch_cooperative_grid=](r   c                   @  s   e Zd Zdd ZdS )z/JITFunction._call_hook.<locals>.JitFunctionInfoc                 S  s   || _ || _|| _d S r6   )rC   r&   jit_function)r/   rC   r&   r6  r3   r3   r4   r%     s   z8JITFunction._call_hook.<locals>.JitFunctionInfo.__init__N)rJ   r;   r   r%   r3   r3   r3   r4   JitFunctionInfo  s    r7  r   )r   devicer   	num_warpsnum_ctas
num_stagesenable_fp_fusionlaunch_cooperative_gridextern_libsconfigsspecialization_data	is_warmupr   )r   r   fncompileis_manual_warmupalready_compiled)rG   
cache_hookcompiled_hookrB  rJ   r;   r  r  paramsr9  r:  r;  r<  r=  r   r>  )r/   r   r   r8  r   r   r?  rA  beforehookr&   rC   	arg_reprsr   r7  r@  r   r3   r3   r4   
_call_hook  s<    8

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/   rJ  r3   r3   r4   add_pre_run_hook  s   zJITFunction.add_pre_run_hookc                 C  sX   ddl m}m}m}m} tj }||}|| _|| _|| _t| j	| j
|}i |||fS )z1
        Precompute as much as possible.
        r   )CompiledKernelrC  	ASTSourcemake_backend)compilerrP  rC  rQ  rR  r   activeget_current_targetr  r   rH  )r/   rP  rC  rQ  rR  r   r  binderr3   r3   r4   create_binder   s   
zJITFunction.create_binderc          !   
     s  | d| jptj dddk|d< tj }tj|}| jD ]	}||i | q | j	| \}}	}
|
|i |\}}t
|t
| }| |d }|d u r|}dd | jD }dd |D }dd	 t||D }d
|vsuJ dd|vs}J dd|vsJ d|D ]}||jvr||vrtd| qt|dd }fdd	|D }dd |D  t dd } fdd	|D }| j||||||g|ddrd S | | |||}| j||	|jd}|||< | j||||||g|dd t }| j D ] \\}}\}}| || }|krtd| d| d| q|ss|d us'J t|r0|}t|}|d }|dkrA|d nd}|d krL|d  nd}|j||g R  } |j|||||j|j| | jj | jj!g	 R   |S )!NdebugTRITON_DEBUG01c                 S     g | ]}|j qS r3   r3  r   r3   r3   r4   rb   $  ro   z#JITFunction.run.<locals>.<listcomp>c                 S  r   r   r3   r   r3   r3   r4   rb   %  r   c                 S  s   i | ]\}}||qS r3   r3   )r`   rN   vr3   r3   r4   r   &  rc   z#JITFunction.run.<locals>.<dictcomp>device_typez=device_type option is deprecated; current target will be usedr8  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  s   |dkS r   r3   )rP   r]   r3   r3   r4   r   /  s    z!JITFunction.run.<locals>.<lambda>c                   s    i | ]}|t t  |qS r3   )r   r   r   )r`   path)
bound_argsr3   r4   r   0  r4  c                 S  r   r   r3   r   r3   r3   r4   rb   2  r   c                 S  s
   t |tS r6   )r#   rM   )rP   r   r3   r3   r4   r   3  s   
 c                   s   i | ]}| t |qS r3   )
parse_attrr   )r`   rN   )attrvalsr  r3   r4   r   4  r   T)rI  )r   r   FrE   z1 has changed since we compiled this kernel, from z to r   r   r   )"r\   rX  osenvironr   rT  get_current_deviceget_current_streamrN  device_cachesrM   parse_optionsrH  r  r   KeyErrorr   rL  rQ  rC  objectr-   r   rI   rM  r   launch_metadatar   r   functionpacked_metadatarP  launch_enter_hooklaunch_exit_hook)!r/   r   r   rp   r   r8  r_  rJ  kernel_cacher   rV  r  r   r   kernelsigkeyssigvalsr   rN   
constexprsr   r0   not_presentr&   rP   r]   globals_dictnewVal	grid_sizegrid_0grid_1grid_2rl  r3   )rc  r  ra  r4   r     sl   $



zJITFunction.runc                 C  s   | j d u r| jS |  |S r6   )_repr_fn_name)r/   rP   r3   r3   r4   r   T  s   zJITFunction.reprc	                 C  sp  |r|ng }|r
|ng }|| _ |j| _|| _t|| _|| _|| _t|d | _	|| _
|j| _|| _g | _t| jj D ]!\}	}
|	|v pK|
j|v }|	|v pT|
j|v }| jt|	|
|| q?tt|}|td|tj d  }| | t| j| _d | _ i | _!d | _"|| _#|| _$dd | jD | _%dd | jD | _&g | _'|j(| _(|j| _|j)| _)|j| _d S )Nr   z^def\s+\w+\s*\(c                 S  r\  r3   r3  r`   pr3   r3   r4   rb     ro   z(JITFunction.__init__.<locals>.<listcomp>c                 S  s   g | ]}|j r|jqS r3   )r   r   r  r3   r3   r4   rb     s    )*rB  r;   rC   versionr=   r   r   r   getsourcelinesstarting_line_numberr}  rJ   r~  rl  rH  	enumerater  r   r&   r   r   textwrapdedent	getsourceresearch	MULTILINEstart_unsafe_update_srcr   rW  rh  hashr-   rr  rX  rF   	arg_namesru  rN  r   __globals__)r/   rB  r  r   r   rX  rF   r   rl  r   r   dnsdns_oar0   r3   r3   r4   r%   W  sB   
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+   r0   )r  r   rJ   r  r0   r_   parser9   rM   r  dictsortedr-   r   )r/   dependencies_finderr3   r3   r4   rL     s   
zJITFunction.cache_keyc                O  s   | j ttj||dd|S )NTr   )r   r  
MockTensor
wrap_dtype)r/   r   rp   r   r3   r3   r4   r     s   zJITFunction.warmupc                   s  ddl m}m} dd l}dd lm  tj }|	|}|d | j
jkr2td|d  d| j
j tt|d }|d } fd	d
t||D }	tt|d }
|d }tt|
|}t|d  }|| ||	|}dd
 |d  D }|d }||d |}|| j| d |< |S )Nr   )rC  rQ  r   r&   zSpecialization data is for z but trying to preload for r   r   c                   s,   i | ]\}}| j |r  |n|qS r3   )r   is_dtyper   tlr3   r4   r     s    z'JITFunction.preload.<locals>.<dictcomp>r   r   r   c                 S  s(   i | ]\}}|t |trt|n|qS r3   )r#   r   r   r   r3   r3   r4   r     s    r   r   )rS  rC  rQ  r   triton.languager   r   rT  rf  loadsrB  rJ   rI   r  r   r  r  r   rh  )r/   r@  rC  rQ  r   r8  deserialized_objr   r   r   r   r   r   r   r0   r   r   rr  r3   r  r4   preload  s4   



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   )rX   r  r0   r#   Moduler   bodyFunctionDef)r/   treer3   r3   r4   r    s
   zJITFunction.parsec                 O  s   t d)Nz:Cannot call @triton.jit'd outside of the scope of a kernel)rI   )r/   rp   r   r3   r3   r4   __call__  s   zJITFunction.__call__c                   s.   |dkrt d| dtt| || d S )Nr0   zCannot set attribute 'zX' directly. Use '_unsafe_update_src()' and manually clear `.hash` of all callersinstead.)AttributeErrorr$   rG   __setattr__)r/   r&   rg   r1   r3   r4   r    s   zJITFunction.__setattr__c                   s   d| _ t d| dS )z
        The only method allowed to modify src.
        Bypasses the __setattr__ restriction by calling super().__setattr__ directly.
        Nr0   )r  r$   r  )r/   new_srcr1   r3   r4   r    s   zJITFunction._unsafe_update_srcc                 C  s   d| j  d| jj dS )NzJITFunction(:r   )rC   rB  rJ   r8   r3   r3   r4   __repr__  s   zJITFunction.__repr__)NNNNNNN)rJ   r;   r   rF  rG  rL  rO  rW  r   r   r%   r   rL   r   r  r  r  r  r  r  r   r3   r3   r1   r4   rG     s(    6G
=
	 rG   rB  JITFunction[T]c                 C  r   r6   r3   )rB  r3   r3   r4   jit     r  r  r   rl  r   r   rX  rF   r   Optional[Callable]rl  r   Optional[Iterable[int]]r   rX  Optional[bool]rF   Callable[[T], JITFunction[T]]c                 C  r   r6   r3   r  r3   r3   r4   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
    rB  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_INTERPRETrZ  r[  r   )InterpretedFunction)r  r   r   rX  rF   r   rl  )rM  rd  getenvinterpreterr  rG   )rB  r  rX  r   r   rl  rF   r   r  r3   r4   	decorator  s"   zjit.<locals>.decoratorNrB  r   r   r  r3   )	rB  r  r   rl  r   r   rX  rF   r  r3   r  r4   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)r2   rJ   r;   r  rl   r3   r3   r4   r  9  s   zMockTensor.wrap_dtypec                 C  s
   || _ d S r6   r   )r/   r   r3   r3   r4   r%   ?     
zMockTensor.__init__c                   C     dS Nr   r3   r3   r3   r3   r4   data_ptrB  r  zMockTensor.data_ptrc                   C  r  r  r3   r3   r3   r3   r4   	ptr_rangeF  r  zMockTensor.ptr_rangeN)	rJ   r;   r   r   staticmethodr  r%   r  r  r3   r3   r3   r4   r  3  s    

r  c                   @  s^   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d ZdS )TensorWrapperc                 C  s*   || _ || _|j| _|j| _| jj| _d S r6   )r   basedatar8  shape)r/   r  r   r3   r3   r4   r%   M  s
   zTensorWrapper.__init__c                 C  r5   r6   )r  r  r8   r3   r3   r4   r  T  r  zTensorWrapper.data_ptrc                 G  s   | j j| S r6   )r  stride)r/   rp   r3   r3   r4   r  W  s   zTensorWrapper.strider   rM   c                 C  s   d| j  d| j dS )NzTensorWrapper[r5  r   )r   r  r8   r3   r3   r4   __str__Z  s   zTensorWrapper.__str__c                 C  r5   r6   )r  element_sizer8   r3   r3   r4   r  ]  r  zTensorWrapper.element_sizec                 C     t | j | jS r6   )r  r  cpur   r8   r3   r3   r4   r  `     zTensorWrapper.cpuc                 C  s   | j |j  d S r6   )r  copy_)r/   otherr3   r3   r4   r  c  r  zTensorWrapper.copy_c                 C  r  r6   )r  r  cloner   r8   r3   r3   r4   r  f  r  zTensorWrapper.clonec                 C     t | j|| jS r6   )r  r  tor   )r/   r8  r3   r3   r4   r  i     zTensorWrapper.toc                 C  r  r6   )r  r  	new_emptyr   )r/   sizesr3   r3   r4   r  l  r  zTensorWrapper.new_emptyNr   rM   )rJ   r;   r   r%   r  r  r  r  r  r  r  r  r  r3   r3   r3   r4   r  K  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   rV   )r   r   r3   r3   r4   reinterpretp  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#   rG   rB  __code__co_filenamer=   r  r  stripr@   )rB  base_fn	file_namelines
begin_lineidxliner3   r3   r4   get_jit_fn_file_line  s   


r  r  )Fr  )r   r  rl  r  r   r  r   r  rX  r  rF   r  r   r  r6   )rB  r  r   r  rl  r  r   r  r   r  rX  r  rF   r  r   r  )9
__future__r   r   rX   r'   r=   rv   rd  r  r  collectionsr   	functoolsr   typingr   r   r   r	   r
   r   r   r   r   r   runtime.driverr   typesr   _utilsr   r   rJ   r   rA   r   NodeVisitorr   r   r   r   r   r   r   r   r   r  r   r   r   r]  rG   r  r  r  r  r  r3   r3   r3   r4   <module>   s    0 
G0
23	

  +<%