o
    m9:j                     @   sD  d dl Z 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de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   Z/home/nk/hobo-godmode/plappi-mvp/.venv/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_tmaM   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_deviceR   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_devicer   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   backendc              
   C   s   |  i }t|dd}|sdS g }t|D ]0\}}tj|rEt|d}|| dt	|
    W d   n1 s@w   Y  qd|S )zReturn a cache key fragment for extern libs (e.g. libdevice.10.bc).

    These files affect codegen but are not covered by triton_key() (Python
    sources only) or backend.hash() (ptxas version and arch only).
    extern_libsN rb-)parse_optionsgetattrsortedospathisfileopenappendhashlibsha256read	hexdigestjoin)rX   optsrY   partsnamera   fr   r   r   _extern_libs_key   s   
$
rn   c                  C   s>   ddl m}  t }|   d|  }t|d  S )Nr   )
triton_keyr\   zutf-8)	%torch._inductor.runtime.triton_compatro   rW   hashre   rf   encoderh   upper)ro   rX   keyr   r   r   triton_hash_with_backend   s   ru   )r	   )	functoolsre   r`   typingr   cacherN   r   r   r   r   r!   r&   r*   r+   r3   r6   	lru_cacher7   rO   rW   strrn   ru   r   r   r   r   <module>   s:    	&

'