
    wi4                     Z   d dl Z d dlZ e j        d          defd            Z e j        d          d             Z e j        d          d             Z e j        d          defd            Z e j        d          d             Z e j        d          d             Z	d	 Z
d
 ZdS )    Nreturnc                  R    	 ddl m}  | d uS # t          $ r Y dS t          $ r Y dS w xY w)Nr   
triton_keyF)triton.compiler.compilerr   ImportErrorRuntimeErrorr   s    c/root/.openclaw/workspace/chatterbox_venv_py311/lib/python3.11/site-packages/torch/utils/_triton.pyhas_triton_packager      s_    777777%%   uu   uus   	 
&	&&c                      t                      radd l} | j                                        rD| j                                        dk    r'| j        j        s	 ddlm}m	} dS # t          $ r Y nw xY wdS )Nr   	   r   )create_1d_tma_descriptorcreate_2d_tma_descriptorTF)r   torchcudais_availableget_device_capabilityversionhip$triton.tools.experimental_descriptorr   r   r   )r   r   r   s      r
   has_triton_tmar            J##%%	
0022f<<M% =       
 t    5   A   
A-,A-c                      t                      radd l} | j                                        rD| j                                        dk    r'| j        j        s	 ddlm}m	} dS # t          $ r Y nw xY wdS )Nr   r   )&experimental_device_tensormap_create1d&experimental_device_tensormap_create2dTF)r   r   r   r   r   r   r   triton.language.extra.cudar   r   r   )r   r   r   s      r
   has_triton_tma_devicer   )   r   r   c                  p    t                      sdS ddlm d } d }d }| ||dfd} |            S )	NFr   )get_interface_for_devicec                 F    | j                                         j        dk    S )N   )Workerget_device_propertiesmajordevice_interfaces    r
   cuda_extra_checkz$has_triton.<locals>.cuda_extra_checkG   s    &<<>>DII    c                 &    dd l }d|j        j        v S )Nr   cpu)triton.backendsbackends)r(   tritons     r
   cpu_extra_checkz#has_triton.<locals>.cpu_extra_checkJ   s    000r*   c                     dS )NT r'   s    r
   _return_truez has_triton.<locals>._return_trueO   s    tr*   )r   xpur,   c                                                       D ]2\  } } |           }|                                r ||          r dS 3dS )NTF)itemsr   )deviceextra_checkr(   r!   triton_supported_devicess      r
    is_device_compatible_with_tritonz4has_triton.<locals>.is_device_compatible_with_tritonX   sg    #;#A#A#C#C 	 	FK77??,,.. ;;?O3P3P ttur*   )r   torch._dynamo.device_interfacer!   )r)   r0   r3   r:   r!   r9   s       @@r
   
has_tritonr<   @   s     uGGGGGGJ J J1 1 1
   !         ,+---r*   c                  b    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?   targets      r
   triton_backendrD   b   sI    555555,,,,,,]--//F<r*   c                     ddl m}  t                      } |              d|                                 }t	          j        |                    d                                                                                    S )Nr   r   -zutf-8)	r   r   rD   hashhashlibsha256encode	hexdigestupper)r   backendkeys      r
   triton_hash_with_backendrO   k   sv    333333GZ\\
,
,GLLNN
,
,C >#**W--..88::@@BBBr*   c                     | j                             d          rd| j         dd          z   }n4| j                             d          rd| j         dd          z   }n| j         }d|z   S )Nfpfloat   bfbfloatztriton.language.)name
startswith)dtypesuffixs     r
   dtype_to_stringrZ   v   sp    zT"" 5:abb>)			t	$	$ EJqrrN*&&r*   c                  2    dd l } d | j        j        _        d S )Nr   c                      t          |           S )N)rZ   )selfs    r
   <lambda>z)patch_triton_dtype_repr.<locals>.<lambda>   s    /$2G2G r*   )r/   languagerX   __repr__)r/   s    r
   patch_triton_dtype_reprra      s$    MMM &H%GFO"""r*   )	functoolsrH   	lru_cacheboolr   r   r   r<   rD   rO   rZ   ra   r2   r*   r
   <module>re      sY        TD     T  , T  , T.D . . . .B T      TC C C' ' 'H H H H Hr*   