a
    zdC                     @  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mZ d dlmZmZmZmZmZmZmZmZ d dlZd dlZd dlmZ zd dlmZ W n ey   dd ZY n0 ed	ZG d
d dejZ e! dd Z"G dd dee Z#G dd de#e Z$ed	ddddZ%eddddddddZ%d#dddddddddZ%G dd  d Z&d!d" Z'dS )$    )annotationsdivisionN)defaultdict
namedtuple)CallableGenericIterableOptionalTypeVarUnioncastoverload)
MockTensor)_cuda_getCurrentRawStreamc                 C  s   t j| jS N)torchZcudaZcurrent_streamZcuda_stream)Zdev_idx r   [/var/www/html/stable-diffusion-webui/venv/lib/python3.9/site-packages/triton/runtime/jit.py<lambda>       r   Tc                      s>   e Zd ZdZdd fddZdd Zdd	 Zd
d Z  ZS )DependenciesFinderz
    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.
    Nonereturnc                   s*   t    t|d | _|| _d S )Nutf-8)super__init__hashlibmd5encode	hexdigestretglobals)selfr#   src	__class__r   r   r   &   s    
zDependenciesFinder.__init__c                 C  s   | j |jd S r   )r#   getid)r$   noder   r   r   
visit_Name+   s    zDependenciesFinder.visit_Namec                 C  sF   |  |j}t|tjr&|  |j}q|d u s6|tu r:d S t||jS r   )visitvalue
isinstanceast	Attributetritongetattrattr)r$   r*   lhsr   r   r   visit_Attribute.   s    z"DependenciesFinder.visit_Attributec                 C  s   |  |j}|d u rd S t|r&d S |jr<|jdr<d S t|tsJJ |jd u rt	
|j}t|j|j}| | |j|_| j|j d| _t| j | _d S )Nztriton.r   )r,   funcinspect	isbuiltin
__module__
startswithr.   JITFunctionhashr/   parser%   r   __globals__r"   r    r   r   r!   )r$   r*   r6   treefinderr   r   r   
visit_Call6   s    


zDependenciesFinder.visit_Call)	__name__r9   __qualname____doc__r   r+   r5   rA   __classcell__r   r   r&   r   r      s
   r   c               	   C  s  dd l } g }ttd(}|t|  g7 }W d    n1 sD0    Y  ttjjd(}|t|  g7 }W d    n1 s0    Y  ttj	j
jd(}|t|  g7 }W d    n1 s0    Y  tjjg tjdR  }| |gD ]V}t|j|jjd(}|t|  g7 }W d    n1 sJ0    Y  q zttddg }W n ty   d}Y n0 dtjd | d d| S )Nr   rblanguageZptxasz	--version -)pkgutilopen__file__r   r   readr!   r1   compilerZ_CZ	libtritonospathjoin__path__iter_modulesmodule_finder	find_specnameorigin
subprocesscheck_output	Exception__version__)rJ   contentsfZlanguage_pathlibZptxas_versionr   r   r   version_keyL   s"    666<
r_   c                   @  s$   e Zd ZU ded< ddddZdS )KernelInterfacer   runr   c                 C  s   t ttjt t| j|dS )z
        A JIT function is launched with: fn[grid](*args, **kwargs).
        Hence JITFunction.__getitem__ returns a callable proxy that
        memorizes the grid.
        )grid)r   r   	functoolspartialr   ra   )r$   rb   r   r   r   __getitem__h   s    zKernelInterface.__getitem__N)rB   r9   rC   __annotations__re   r   r   r   r   r`   e   s   
r`   c                      s   e Zd ZdZdZedd Zedd Zdd Ze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 fddZdd  Z  ZS )"r;   N   c                 C  s   t | dr| jS t| trdS t| trjd| kr<| dkr<dS d| krP| dkrPdS d	| krd| d
krddS dS n2t| trxdS | d u rd S tdt|  d|  d S )Ndtypei1i   ii32l        l    u32l            l    u64i64fp32zUnsupported type z for )hasattrrh   r.   boolintfloat	TypeErrortypeargr   r   r   _key_ofw   s"    



zJITFunction._key_ofc                 C  sD   t | dr|  tj dkS t| tr:| d dk| dkfS | d u fS )Ndata_ptrr   rg      ro   rx   r;   divisibilityr.   rq   ru   r   r   r   _spec_of   s
    

zJITFunction._spec_ofc                   sR   dd   fddt |D }fddt |D }tdddgt|t|S )	Nc                 S  sD   t | dr|  tj dkS t| tr4| tj dkS | d u r@dS dS )Nrx   r   TFrz   )xr   r   r   is_divisible_by_16   s    

z3JITFunction._get_config.<locals>.is_divisible_by_16c                   s&   h | ]\}} |r|j vr|qS r   )do_not_specialize.0irv   r~   r$   r   r   	<setcomp>   r   z*JITFunction._get_config.<locals>.<setcomp>c                   s0   h | ](\}}t |tr|d kr| jvr|qS )ry   )r.   rq   r   r   r$   r   r   r      r   Zinstance_descriptordivisible_by_16
equal_to_1)	enumerater   tuple)r$   argsr   r   r   r   r   _get_config   s    zJITFunction._get_configc              $   C  s   t | tjtjjfrtjdtjdtjdtjdtj	dtj
dtjdtjdtjd	tjd
tjj
dtjjdtjjdtjjdtjjdtjjdtjjdtjjdi|  }d| S | d u rdS t | tsJ | S )Nri   Zfp16Zbf16rn   Zfp64u8i8Zi16rj   rm   u16rk   rl   Zfp8*z*i8)r.   r   rh   r1   rG   rp   float16Zbfloat16float32float64Zuint8Zint8int16int32int64Zuint16Zuint32Zuint64Zfloat8str)keytyr   r   r   _type_of   s6    
zJITFunction._type_ofc                   s    d  fddt|D }|S )N,c                   s   g | ]\}}  |qS r   )r   )r   r   kr   r   r   
<listcomp>   r   z/JITFunction._make_signature.<locals>.<listcomp>)rQ   r   )r$   Zsig_key	signaturer   r   r   _make_signature   s    zJITFunction._make_signaturec                 C  s   t t| j|}|S r   )dictzip
constexprs)r$   Zconstexpr_key	constantsr   r   r   _make_constants   s    zJITFunction._make_constantsc	              	   C  s   t jd u rdS | jj}	| jj}
ddd t| j|d D }|	 d| d| d| d	}t|}G d
d d}t	|||||||d}t j||||
|	d|i|dddS )NF, c                 S  s   g | ]\}}| d | qS )z: r   r   rV   r   r   r   r   r      r   z*JITFunction._call_hook.<locals>.<listcomp>ry   z[num_warps=z, num_stages=]()c                   @  s   e Zd Zdd ZdS )z.JITFunction._call_hook.<locals>.LegacyCompilerc                 S  s   || _ || _d S r   )modulerV   )r$   r   rV   r   r   r   r      s    z7JITFunction._call_hook.<locals>.LegacyCompiler.__init__N)rB   r9   rC   r   r   r   r   r   LegacyCompiler   s   r   )r   devicer   	num_warps
num_stagesextern_libsconfigsr   )r   reprfncompileZis_manual_warmupZalready_compiled)
r;   
cache_hookr   rB   r9   rQ   r   	arg_namesr   r   )r$   r   r   r   r   r   r   r   r   rV   r   Z	arg_reprsr   r   kwargsr   r   r   
_call_hook   s    
 zJITFunction._call_hookc                   s   fddt  jD } fddt  jD }d|}ddd |D }d|}g }t |D ]N\}}| jv rxqd|d| dtj d| d	| d
tj d| d| dg7 }qdd|}	ddd  jD }
d jj dd j d| dt|dkr| dnd dt|	dkr*|	 dnd d|
 d| d| dddd  jD  d}t	 t
  j j jttd}t|| | jj S )Nc                   s    g | ]\}}| j vr| qS r   r   r   r   r   r   r      r   z.JITFunction._make_launcher.<locals>.<listcomp>c                   s    g | ]\}}| j v r| qS r   r   r   r   r   r   r      r   r   c                 S  s   g | ]}d | dqS )z_key_of(r   r   r   rv   r   r   r   r      r   (z.data_ptr() % z == 0) if hasattr(z, "data_ptr") else (z % z == 0, z == 1) if isinstance(z, int) else (False,)r   c                 S  s   g | ]}d | d| qS )"z": r   r   r   r   r   r      r   z
def z`, grid, num_warps=4, num_stages=3, extern_libs=None, stream=None, warmup=False):
    sig_key =  z,
    constexpr_key = r   r   z
    spec_key = a  
    key = (version_key, sig_key, constexpr_key, spec_key)
    if not extern_libs is None:
      key = (key, tuple(extern_libs.items()))
    assert num_warps > 0 and (num_warps & (num_warps - 1)) == 0, "num_warps must be a power of 2"
    if callable(grid):
        grid = grid({a-  })
    grid_size = len(grid)
    grid_0 = grid[0]
    grid_1 = grid[1] if grid_size > 1 else 1
    grid_2 = grid[2] if grid_size > 2 else 1
    device = torch.cuda.current_device()
    torch.cuda.set_device(device)
    if stream is None and not warmup:
      stream = get_cuda_stream(device)
    try:
      bin = cache[device][key]
      if not warmup:
          bin.c_wrapper(grid_0, grid_1, grid_2, bin.num_warps, bin.shared, stream, bin.cu_function, triton.compiler.CompiledKernel.launch_enter_hook, triton.compiler.CompiledKernel.launch_exit_hook, bin, z)
      return bin
    # kernel not cached -- compile
    except KeyError:
      # build dict of constant values
      args = [z]
      all_args = c                 S  s   g | ]
}| qS r   r   r   r   r   r   r     r   a  ,
      configs = self._get_config(*all_args),
      constants = self._make_constants(constexpr_key)
      constants.update({i: None for i, arg in enumerate(all_args) if arg is None})
      constants.update({i: 1 for i in configs[0].equal_to_1})
      # build kernel signature -- doesn't include specialized arguments
      signature = { i: self._type_of(_key_of(arg)) for i, arg in enumerate(all_args) if i not in self.constexprs }
      # build stub signature -- includes arguments that are specialized
      for i, arg in constants.items():
        if callable(arg):
          raise TypeError(f"Callable constexpr at index {i} is not supported")
      if not self._call_hook(key, signature, device, constants, num_warps, num_stages, extern_libs, configs):
        bin = triton.compile(self, signature=signature, device=device, constants=constants, num_warps=num_warps, num_stages=num_stages, extern_libs=extern_libs, configs=configs)
        if not warmup:
            bin.c_wrapper(grid_0, grid_1, grid_2, bin.num_warps, bin.shared, stream, bin.cu_function, triton.compiler.CompiledKernel.launch_enter_hook, triton.compiler.CompiledKernel.launch_exit_hook, bin, *args)
        self.cache[device][key] = bin
        return bin
      return None
)r_   get_cuda_streamr$   r|   rw   cacher1   r   )r   r   rQ   r   r;   r{   r   rB   lenr_   r   r|   rw   r   r1   r   exec)r$   Zregular_argsZconstexpr_argsr   Zsig_keysZconstexpr_keysZspecializationsr   rv   Z	spec_keysZ	grid_argsr%   Zscoper   r   r   _make_launcher   sZ    





/

zJITFunction._make_launcherc                   s&  | _ |j _| _t|}dd |j D  _t	dd |j D  _
|d u rZg n| _ fdd jD  _tt| _ j jdd   _tt _d  _g  _d  _ fdd	|j D  _|j _ fd
d j D  _   _|j _|j _|j  _ |j _d S )Nc                 S  s   g | ]
}|j qS r   )rV   r   vr   r   r   r   ,  r   z(JITFunction.__init__.<locals>.<listcomp>c                 s  s   | ]}|j tjkV  qd S r   )defaultr7   _emptyr   r   r   r   	<genexpr>-  r   z'JITFunction.__init__.<locals>.<genexpr>c                   s&   h | ]}t |tr j|n|qS r   )r.   r   r   indexr   r   r   r   r   0  r   z'JITFunction.__init__.<locals>.<setcomp>defc                   s   i | ]\}} j ||qS r   r   r   r   r   r   r   
<dictcomp><  r   z(JITFunction.__init__.<locals>.<dictcomp>c                   s   g | ]} j |qS r   r   )r   annr   r   r   r   ?  r   )!r   r9   r   versionr7   r   
parametersvaluesr   anyZhas_defaultsr   textwrapdedent	getsourcer%   findr   r   r   r<   kernel_decoratorskernelrf   itemsr   keysr   r   ra   rD   rB   r>   )r$   r   r   r   r   r   r   r   r   &  s,    


zJITFunction.__init__c                 C  s<   | j d u r6t| j| jd}||   |jt  | _ | j S )N)r#   r%   )r<   r   r>   r%   r,   r=   r"   r_   )r$   Zdependencies_finderr   r   r   	cache_keyH  s
    
zJITFunction.cache_keyc                 O  s    | j ttj|i |ddiS )NwarmupT)ra   mapr   Z
wrap_dtyper$   r   r   r   r   r   r   Q  s    zJITFunction.warmupc                 C  sH   t | j}t|t jsJ t|jdks.J t|jd t jsDJ |S )Nry   r   )r/   r=   r%   r.   Moduler   bodyFunctionDef)r$   r?   r   r   r   r=   W  s
    zJITFunction.parsec                 O  s   t dd S )Nz:Cannot call @triton.jit'd outside of the scope of a kernel)RuntimeErrorr   r   r   r   __call__^  s    zJITFunction.__call__c                   s2   |dkrd | _ tt| || |dkr.d | _d S )Nr   r%   )r   r   r;   __setattr__r<   )r$   rV   r-   r&   r   r   r   a  s
    zJITFunction.__setattr__c                 C  s   d| j  d| jj dS )NzJITFunction(:r   )r   r   rB   r   r   r   r   __repr__l  s    zJITFunction.__repr__)NN)rB   r9   rC   r   r{   staticmethodrw   r|   r   r   r   r   r   r   r   propertyr   r   r=   r   r   r   rE   r   r   r&   r   r;   q   s*   


H
"
r;   JITFunction[T]r   r   c                 C  s   d S r   r   r   r   r   r   jitu  s    r   r   r   zOptional[Iterable[int]]zCallable[[T], JITFunction[T]])r   r   c                 C  s   d S r   r   r   r   r   r   r   z  s    zOptional[T]z4Union[JITFunction[T], Callable[[T], JITFunction[T]]])r   r   r   c                  s.   ddd fdd}| dur&|| S |S dS )a+  
    Decorator for JIT-compiling a function using the Triton compiler.

    :note: When a jit'd function is called, :code:`torch.tensor` arguments are
        implicitly converted to pointers using the :code:`.data_ptr()` method.

    :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   c                   s   t | sJ t|  dS )Nr   )callabler;   r   r   r   r   r   	decorator  s    zjit.<locals>.decoratorNr   )r   r   r   r   r   r   r   r     s    c                   @  s*   e Zd Zdd Zdd ZddddZd	S )
TensorWrapperc                 C  s    || _ || _|j| _|j| _d S r   )rh   baseZis_cudar   )r$   r   rh   r   r   r   r     s    zTensorWrapper.__init__c                 C  s
   | j  S r   )r   rx   r   r   r   r   rx     s    zTensorWrapper.data_ptrr   r   c                 C  s   d| j  d| j dS )NzTensorWrapper[r   r   )rh   r   r   r   r   r   __str__  s    zTensorWrapper.__str__N)rB   r9   rC   r   rx   r   r   r   r   r   r     s   r   c                 C  sX   t | tr*|| jjkr| jS t| j|S n*t | tjr@t| |S tdt|  dd S )NzCannot reinterpret a .)r.   r   r   rh   r   ZTensorrs   rt   )Ztensorrh   r   r   r   reinterpret  s    

r   )N)(
__future__r   r   r/   rc   r   r7   rO   rX   r   collectionsr   r   typingr   r   r   r	   r
   r   r   r   r   r1   Ztriton.utilsr   Ztorch._Cr   r   ImportErrorr   NodeVisitorr   	lru_cacher_   r`   r;   r   r   r   r   r   r   r   <module>   sJ   (-
  	 &