
    i                     >   d dl Z d dlZd dlmZ e j                  defd       Ze j                  ddeeef   deeef   fd       Z	e j                  defd       Z
e j                  defd       Ze j                  defd	       Ze j                  defd
       Ze j                  defd       Ze j                  defd       Z e j                   d      defd       Ze j                  defd       Ze j                  defd       Ze j                  defd       Zy)    N)Anyreturnc                  ,    	 dd l } y# t        $ r Y yw xY w)Nr   TF)tritonImportError)r   s    M/var/www/html/engine/venv/lib/python3.12/site-packages/torch/utils/_triton.pyhas_triton_packager	      s     s    	fallbackc                     	 dd l }t        d |j                  j                  d      d d D              \  }}||fS # t        $ r | cY S w xY w)Nr   c              3   2   K   | ]  }t        |        y wN)int).0vs     r   	<genexpr>z%get_triton_version.<locals>.<genexpr>   s     OSVOs   .   )r   tuple__version__splitr   )r
   r   majorminors       r   get_triton_versionr      sT    OV-?-?-E-Ec-J2A-NOOuu~ s   8; A	A	c                      dd l } | j                  j                         xr6 | j                  j                         dk\  xr | j                  j
                   S )Nr   	   r   )torchcudais_availableget_device_capabilityversionhipr   s    r   _device_supports_tmar$      sK     	

! 	"JJ,,.&8	"!!!    c                      t               r!t               r	 ddlm} m} 	 ddlm}  |       S y# t        $ r Y yw xY w# t        $ r Y yw xY w)Nr   )create_1d_tma_descriptorcreate_2d_tma_descriptor)enable_in_pytorchTF)r	   r$   $triton.tools.experimental_descriptorr'   r(   r)   r   )r'   r(   r)   s      r    has_triton_experimental_host_tmar+   &   sV    !
 V,..  #    s#   < - 	9< 9< 	AAc                  Z    t               rt               r	 ddlm}  yy# t        $ r Y yw xY w)Nr   TensorDescriptorTF)r	   r$   triton.tools.tensor_descriptorr.   r   r-   s    r   %has_triton_tensor_descriptor_host_tmar0   <   s9    !    s    	**c                  .    t               xs
 t               S r   )r0   r+    r%   r   has_triton_tmar3   L   s    02X6V6XXr%   c                  N   t               rudd l} | j                  j                         r3| j                  j	                         dk\  r| j
                  j                  r| j                  j                         r
	 ddlm	}m
} yy# t        $ r Y nw xY w	 ddlm} y# t        $ r Y yw xY w)Nr   r   )&experimental_device_tensormap_create1d&experimental_device_tensormap_create2dTmake_tensor_descriptorF)r	   r   r   r   r    r!   r"   xputriton.language.extra.cudar5   r6   r   triton.languager8   )r   r5   r6   r8   s       r   has_triton_tma_devicer<   Q   s     JJ##%

002f<MM%%YY##%
    B s$   7B 	BBB 	B$#B$c                     dd l } | j                  j                         rf| j                  j                         dk\  rI| j                  j                         dk  r,| j                  j
                  st               xr
 t               S y)Nr   )
   r   )   r   F)r   r   r   r    r!   r"   r<   r0   r#   s    r   #has_datacenter_blackwell_tma_devicer@   q   sc     	

!JJ,,.'9JJ,,.8!!$&R+P+RRr%   c                     t               rsdd l} | j                  j                         r3| j                  j	                         dk\  r| j
                  j                  r| j                  j                         r	 ddlm	} yy# t        $ r Y yw xY w)Nr   r   r7   TF)r	   r   r   r   r    r!   r"   r9   r;   r8   r   )r   r8   s     r   has_triton_stable_tma_apirB      sv     JJ##%

002f<MM%%YY##%B   s   7A? ?	B
Bc                      t               syddlm}  | ryddlm dt
        dt        fd}dt
        dt        fd}dt
        dt        fd	}||||d
dt        ffd} |       S )NFr   )triton_disable_device_detection)get_interface_for_devicedevice_interfacer   c                 P    | j                   j                         j                  dk\  S )N   )Workerget_device_propertiesr   rF   s    r   cuda_extra_checkz$has_triton.<locals>.cuda_extra_check   s"    &&<<>DDIIr%   c                 :    dd l }d|j                  j                  v S )Nr   cpu)triton.backendsbackends)rF   r   s     r   cpu_extra_checkz#has_triton.<locals>.cpu_extra_check   s    0000r%   c                      y)NTr2   rK   s    r   _return_truez has_triton.<locals>._return_true   s    r%   )r   r9   rN   mtiac                  |    j                         D ](  \  } } |       }|j                         s ||      s( y y)NTF)itemsr   )deviceextra_checkrF   rE   triton_supported_devicess      r    is_device_compatible_with_tritonz4has_triton.<locals>.is_device_compatible_with_triton   sG    #;#A#A#C 	FK7?,,.;?O3P	 r%   )r	   torch._inductor.configrD   torch._dynamo.device_interfacerE   r   bool)rD   rL   rQ   rS   rZ   rE   rY   s        @@r   
has_tritonr^      s    F&GJ3 J4 J1# 1$ 1
s t  !	 d  ,--r%   c                  ^    ddl m}  ddlm} |j                  j                         } | |      S )Nr   )make_backend)driver)triton.compiler.compilerr`   triton.runtime.driverra   activeget_current_target)r`   ra   targets      r   triton_backendrg      s%    5,]]--/Fr%   c                      ddl m}  t               } |         d|j                          }t	        j
                  |j                  d            j                         j                         S )Nr   )
triton_key-zutf-8)	%torch._inductor.runtime.triton_compatri   rg   hashhashlibsha256encode	hexdigestupper)ri   backendkeys      r   triton_hash_with_backendrt      sR    @G\N!GLLN+
,C >>#**W-.88:@@BBr%   ))r   r   )	functoolsrm   typingr   cacher]   r	   r   r   r   r$   r+   r0   r3   r<   r@   	lru_cacherB   r^   rg   strrt   r2   r%   r   <module>rz      s      D   sCx eCHo   d   $  * t   Y Y Y t  > T   T4  $ $.D $. $.N       C# C Cr%   