
    Vh4                     L   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y)    Nreturnc                  J    	 ddl m}  | d uS # t        $ r Y yt        $ r Y yw xY w)Nr   
triton_keyF)triton.compiler.compilerr   ImportErrorRuntimeErrorr   s    C/home/dcms/DCMS/lib/python3.12/site-packages/torch/utils/_triton.pyhas_triton_packager      s2    7%%  s   	 	"""c                      t               r[dd l} | j                  j                         r=| j                  j	                         dk\  r | j
                  j                  s
	 ddlm}m	} yy# t        $ r Y yw xY w)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     JJ##%

002f<MM%%
       A' '	A32A3c                      t               r[dd l} | j                  j                         r=| j                  j	                         dk\  r | j
                  j                  s
	 ddlm}m	} yy# t        $ r Y yw xY w)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                  `    t               syddlm d } d }d }| ||dfd} |       S )	NFr   )get_interface_for_devicec                 P    | j                   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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00r*   c                      y)NT r'   s    r
   _return_truez has_triton.<locals>._return_trueO   s    r*   )r   xpur,   c                  |    j                         D ](  \  } } |       }|j                         s ||      s( y y)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 	FK7?,,.;?O3P	 r*   )r   torch._dynamo.device_interfacer!   )r)   r0   r3   r:   r!   r9   s       @@r
   
has_tritonr<   @   sB    GJ1
 !  ,--r*   c                  ^    ddl m}  ddlm} |j                  j                         } | |      S )Nr   )make_backend)driver)r   r>   triton.runtime.driverr?   activeget_current_target)r>   r?   targets      r
   triton_backendrD   b   s%    5,]]--/Fr*   c                      ddl m}  t               } |         d|j                          }t	        j
                  |j                  d            j                         j                         S )Nr   r   -zutf-8)	r   r   rD   hashhashlibsha256encode	hexdigestupper)r   backendkeys      r
   triton_hash_with_backendrO   k   sR    3G\N!GLLN+
,C >>#**W-.88:@@BBr*   c                     | j                   j                  d      rd| j                   dd  z   }d|z   S | j                   j                  d      rd| j                   dd  z   }d|z   S | j                   }d|z   S )Nfpfloat   bfbfloatztriton.language.)name
startswith)dtypesuffixs     r
   dtype_to_stringrZ   v   s    zzT"5::ab>)
 &&	 
		t	$EJJqrN* && &&r*   c                  D    dd l } d | j                  j                  _        y )Nr   c                     t        |       S )N)rZ   )selfs    r
   <lambda>z)patch_triton_dtype_repr.<locals>.<lambda>   s    /$2G r*   )r/   languagerX   __repr__)r/   s    r
   patch_triton_dtype_reprra      s     &HFOO"r*   )	functoolsrH   	lru_cacheboolr   r   r   r<   rD   rO   rZ   ra   r2   r*   r
   <module>re      s      TD   T , T , T.D . .B T    TC C'Hr*   