a
    h                     @   s   d dl Z d dlZd dlmZ e jedddZe jedddZe jeddd	Ze jedd
dZ	e jedddZ
e jedddZe dedddZe jedddZe jedddZe jedddZdS )    N)Anyreturnc                  C   sB   zddl m}  | d uW S  ty*   Y dS  ty<   Y dS 0 d S )Nr   
triton_keyF)triton.compiler.compilerr   ImportErrorRuntimeErrorr    r
   A/var/www/auris/lib/python3.9/site-packages/torch/utils/_triton.pyhas_triton_package   s    
r   c                  C   s*   dd l } | j o(| j dko(| jj S )Nr   	   r   )torchcudais_availableget_device_capabilityversionhip)r   r
   r
   r   _device_supports_tma   s    
r   c                  C   s:   t  r6t r6zddlm} m} W dS  ty4   Y n0 dS )Nr   create_1d_tma_descriptorcreate_2d_tma_descriptorTF)r   r   Z$triton.tools.experimental_descriptorr   r   r   r   r
   r
   r    has_triton_experimental_host_tma   s    r   c                  C   s6   t  r2t r2zddlm}  W dS  ty0   Y n0 dS )Nr   TensorDescriptorTF)r   r   Ztriton.tools.tensor_descriptorr   r   r   r
   r
   r   %has_triton_tensor_descriptor_host_tma.   s    r   c                   C   s   t  p
t S )N)r   r   r
   r
   r
   r   has_triton_tma>   s    r   c                  C   s   t  r~dd l} | j r~| j dkr~| jjs~zddlm}m	} W dS  t
yV   Y n0 zddlm} W dS  t
y|   Y n0 dS )Nr   r   )&experimental_device_tensormap_create1d&experimental_device_tensormap_create2dTmake_tensor_descriptorF)r   r   r   r   r   r   r   Ztriton.language.extra.cudar   r   r   triton.languager!   )r   r   r   r!   r
   r
   r   has_triton_tma_deviceC   s&    r#   c                  C   sX   t  rTdd l} | j rT| j dkrT| jjsTzddlm} W dS  t	yR   Y n0 dS )Nr   r   r    TF)
r   r   r   r   r   r   r   r"   r!   r   )r   r!   r
   r
   r   has_triton_stable_tma_apic   s    r$   c                     sl   t  s
dS ddlm  ttddd} ttddd}ttdd	d
}| ||dtd fdd}| S )NFr   )get_interface_for_device)device_interfacer   c                 S   s   | j  jdkS )N   )ZWorkerZget_device_propertiesmajorr&   r
   r
   r   cuda_extra_check}   s    z$has_triton.<locals>.cuda_extra_checkc                 S   s   dd l }d|jjv S )Nr   cpu)Ztriton.backendsbackends)r&   Ztritonr
   r
   r   cpu_extra_check   s    z#has_triton.<locals>.cpu_extra_checkc                 S   s   dS )NTr
   r)   r
   r
   r   _return_true   s    z has_triton.<locals>._return_true)r   Zxpur+   r   c                     s4     D ]&\} } | }| r||r dS qdS )NTF)itemsr   )ZdeviceZextra_checkr&   r%   Ztriton_supported_devicesr
   r    is_device_compatible_with_triton   s
    z4has_triton.<locals>.is_device_compatible_with_triton)r   Ztorch._dynamo.device_interfacer%   r   bool)r*   r-   r.   r1   r
   r0   r   
has_tritonv   s    r3   c                  C   s*   ddl m}  ddlm} |j }| |S )Nr   )make_backend)driver)r   r4   Ztriton.runtime.driverr5   activeZget_current_target)r4   r5   targetr
   r
   r   triton_backend   s    
r8   c                  C   s>   ddl m}  t }|   d|  }t|d  S )Nr   r   -zutf-8)	r   r   r8   hashhashlibsha256encode	hexdigestupper)r   backendkeyr
   r
   r   triton_hash_with_backend   s    rB   )	functoolsr;   typingr   cacher2   r   r   r   r   r   r#   	lru_cacher$   r3   r8   strrB   r
   r
   r
   r   <module>   s,   
!