o
    h                     @   s   d dl Z d dlZd dlmZ e jdefddZe jdefddZe jdefdd	Ze jdefd
dZ	e jdefddZ
e jdefddZe ddefddZe jdefddZe jdefddZe jdefddZdS )    N)Anyreturnc                  C   s>   zddl m}  | d uW S  ty   Y dS  ty   Y dS w )Nr   
triton_keyF)triton.compiler.compilerr   ImportErrorRuntimeErrorr    r	   Y/home/www/facesmatcher.com/frenv_anti/lib/python3.10/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  rt rzddlm} m} W dS  ty   Y dS w 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   s8   t  rt rz	ddlm}  W dS  ty   Y dS w 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
y+   Y nw z	ddlm} W dS  t
y?   Y dS w 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   sZ   t  r+dd l} | j r+| j dkr+| jjs+z	ddlm} W dS  t	y*   Y dS w 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                     sr   t  sdS ddlm  dtdtfdd} dtdtfdd	}dtdtfd
d}| ||ddtf 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*   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-   r0   r	   r/   r
   
has_tritonv   s   r2   c                  C   s*   ddl m}  ddlm} |j }| |S )Nr   )make_backend)driver)r   r3   Ztriton.runtime.driverr4   ZactiveZget_current_target)r3   r4   targetr	   r	   r
   triton_backend   s   
r6   c                  C   s>   ddl m}  t }|   d|  }t|d  S )Nr   r   -zutf-8)	r   r   r6   hashhashlibsha256encode	hexdigestupper)r   backendkeyr	   r	   r
   triton_hash_with_backend   s   r@   )	functoolsr9   typingr   cacher1   r   r   r   r   r   r"   	lru_cacher#   r2   r6   strr@   r	   r	   r	   r
   <module>   s.    
!