o
    i                     @   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deeef deeef 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 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 } W dS  ty   Y dS w )Nr   TF)tritonImportError)r    r   U/sda-disk/www/egybert/egybert_env/lib/python3.10/site-packages/torch/utils/_triton.pyhas_triton_package   s   r   r   r   fallbackc                 C   sP   zdd l }tdd |jdd d D \}}||fW S  ty'   |  Y S w )Nr   c                 s   s    | ]}t |V  qd S N)int).0vr   r   r   	<genexpr>   s    z%get_triton_version.<locals>.<genexpr>.   )r   tuple__version__splitr   )r
   r   majorminorr   r   r   get_triton_version   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   sd   t  r0t r0zddlm} m} zddlm} | W W S  ty%   Y W dS w  ty/   Y dS w dS )Nr   )create_1d_tma_descriptorcreate_2d_tma_descriptor)enable_in_pytorchTF)r   r!   $triton.tools.experimental_descriptorr"   r#   r$   r   )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 r   )r*   r&   r   r   r   r   has_triton_tmaL   s   r+   c                  C   s   t  rEdd l} | j r| j dkr| jjr| j rEzddlm	}m
} W dS  ty0   Y nw z	ddlm} W dS  tyD   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   xputriton.language.extra.cudar,   r-   r   triton.languager/   )r   r,   r-   r/   r   r   r   has_triton_tma_deviceQ   s.   r3   c                  C   sF   dd l } | j r!| j dkr!| j dk r!| jjs!t o t S dS )Nr   )
   r   )   r   F)r   r   r   r   r   r   r3   r*   r    r   r   r   #has_datacenter_blackwell_tma_deviceq   s   r6   c                  C   sd   t  r0dd l} | j r| j dkr| jjr| j r0z	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   r0   r2   r/   r   )r   r/   r   r   r   has_triton_stable_tma_api   s"   r7   c                     s   t  sdS ddlm}  | r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   )triton_disable_device_detection)get_interface_for_devicedevice_interfacer   c                 S   s   | j  jdkS )N   )Workerget_device_propertiesr   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:   r   r   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   r0   r@   mtiac                     s4     D ]\} } | }| r||r dS qdS )NTF)itemsr   )deviceextra_checkr:   r9   triton_supported_devicesr   r    is_device_compatible_with_triton   s   z4has_triton.<locals>.is_device_compatible_with_triton)r   torch._inductor.configr8   torch._dynamo.device_interfacer9   r   bool)r8   r?   rC   rD   rK   r   rI   r   
has_triton   s    rO   c                  C   s*   ddl m}  ddlm} |j }| |S )Nr   )make_backend)driver)triton.compiler.compilerrP   triton.runtime.driverrQ   activeget_current_target)rP   rQ   targetr   r   r   triton_backend   s   
rW   c                  C   s>   ddl m}  t }|   d|  }t|d  S )Nr   )
triton_key-zutf-8)	%torch._inductor.runtime.triton_compatrX   rW   hashhashlibsha256encode	hexdigestupper)rX   backendkeyr   r   r   triton_hash_with_backend   s   rc   )r	   )	functoolsr\   typingr   cacherN   r   r   r   r   r!   r&   r*   r+   r3   r6   	lru_cacher7   rO   rW   strrc   r   r   r   r   <module>   s6    	&

'