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	   N/var/www/html/scripts/venv/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   $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   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   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   )Worker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)triton.backendsbackends)r(   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   xpur/   c                     s4     D ]\} } | }| r||r dS qdS )NTF)itemsr   )deviceextra_checkr(   r'   triton_supported_devicesr	   r
    is_device_compatible_with_triton   s   z4has_triton.<locals>.is_device_compatible_with_triton)r   torch._dynamo.device_interfacer'   r   bool)r.   r3   r4   r;   r	   r9   r
   
has_tritonv   s   r>   c                  C   s*   ddl m}  ddlm} |j }| |S )Nr   )make_backend)driver)r   r?   triton.runtime.driverr@   activeget_current_target)r?   r@   targetr	   r	   r
   triton_backend   s   
rE   c                  C   s>   ddl m}  t }|   d|  }t|d  S )Nr   r   -zutf-8)	r   r   rE   hashhashlibsha256encode	hexdigestupper)r   backendkeyr	   r	   r
   triton_hash_with_backend   s   rO   )	functoolsrH   typingr   cacher=   r   r   r   r   r   r%   	lru_cacher&   r>   rE   strrO   r	   r	   r	   r
   <module>   s.    
!