
    [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dS )    N)Anyreturnc                  2    	 dd l } dS # t          $ r Y dS w xY w)Nr   TF)tritonImportError)r   s    F/var/www/icac/venv/lib/python3.11/site-packages/torch/utils/_triton.pyhas_triton_packager	      s7    t   uus    
r   r   fallbackc                     	 dd l }t          d |j                            d          d d         D                       \  }}||fS # t          $ r | cY S w xY w)Nr   c              3   4   K   | ]}t          |          V  d S N)int).0vs     r   	<genexpr>z%get_triton_version.<locals>.<genexpr>   s(      OOSVVOOOOOO    .   )r   tuple__version__splitr   )r   r   majorminors       r   get_triton_versionr      sx    OOV-?-E-Ec-J-J2A2-NOOOOOuu~   s   AA AAc                      dd l } | j                                        o)| j                                        dk    o| j        j         S )Nr   	   r   )torchcudais_availableget_device_capabilityversionhipr   s    r   _device_supports_tmar&      sN    LLL 	
!! 	"J,,..&8	"!!r   c                      t                      rIt                      r;	 ddlm} m} 	 ddlm}  |            S # t          $ r Y dS w xY w# t          $ r Y nw xY 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*   s      r    has_triton_experimental_host_tmar,   &   s     !! 	       
 VVVVVV,,..."      44     5s,   A 7 
AA AA 
AAc                  p    t                      r't                      r	 ddlm}  dS # t          $ r Y nw xY wdS )Nr   TensorDescriptorTF)r	   r&   triton.tools.tensor_descriptorr/   r   r.   s    r   %has_triton_tensor_descriptor_host_tmar1   <   so     	!! 	      t    5s   & 
33c                  :    t                      pt                      S r   )r1   r,    r   r   has_triton_tmar4   L   s    022X6V6X6XXr   c                  H   t                      rdd l} | j                                        r)| j                                        dk    r| j        j        r| j                                        r4	 ddlm	}m
} dS # t          $ r Y nw xY w	 ddlm} dS # t          $ r Y nw xY 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.cudar6   r7   r   triton.languager9   )r   r6   r7   r9   s       r   has_triton_tma_devicer=   Q   s     J##%%	
0022f<<M% =Y##%% =       
 t   BBBBBBt    5s$   /A9 9
BB
B 
BBc                     dd l } | j                                        rb| j                                        dk    rE| j                                        dk     r(| j        j        st                      ot                      S dS )Nr   )
   r   )   r   F)r   r    r!   r"   r#   r$   r=   r1   r%   s    r   #has_datacenter_blackwell_tma_devicerA   q   s    LLL 	
!!SJ,,..'99J,,..88! 9 %&&R+P+R+RR5r   c                     t                      rxdd l} | j                                        r)| j                                        dk    r| j        j        r| j                                        r	 ddlm	} dS # t          $ r Y nw xY wdS )Nr   r   r8   TF)r	   r   r    r!   r"   r#   r$   r:   r<   r9   r   )r   r9   s     r   has_triton_stable_tma_apirC      s      J##%%
	
0022f<<M% =Y##%% =BBBBBBt   5s   /A7 7
BBc                      t                      sdS ddlm}  | rdS 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                 F    | j                                         j        dk    S )N   )Workerget_device_propertiesr   rG   s    r   cuda_extra_checkz$has_triton.<locals>.cuda_extra_check   s    &<<>>DIIr   c                 &    dd l }d|j        j        v S )Nr   cpu)triton.backendsbackends)rG   r   s     r   cpu_extra_checkz#has_triton.<locals>.cpu_extra_check   s    000r   c                     dS )NTr3   rL   s    r   _return_truez has_triton.<locals>._return_true   s    tr   )r    r:   rO   mtiac                                                       D ]2\  } } |           }|                                r ||          r dS 3dS )NTF)itemsr!   )deviceextra_checkrG   rF   triton_supported_devicess      r    is_device_compatible_with_tritonz4has_triton.<locals>.is_device_compatible_with_triton   sg    #;#A#A#C#C 	 	FK77??,,.. ;;?O3P3P ttur   )r	   torch._inductor.configrE   torch._dynamo.device_interfacerF   r   bool)rE   rM   rR   rT   r[   rF   rZ   s        @@r   
has_tritonr_      s
    uFFFFFF& uGGGGGGJ3 J4 J J J J1# 1$ 1 1 1 1
s t     !	   d        ,+---r   c                  b    ddl m}  ddlm} |j                                        } | |          S )Nr   )make_backend)driver)triton.compiler.compilerra   triton.runtime.driverrb   activeget_current_target)ra   rb   targets      r   triton_backendrh      sI    555555,,,,,,]--//F<r   c                     ddl m}  t                      } |              d|                                 }t	          j        |                    d                                                                                    S )Nr   )
triton_key-zutf-8)	%torch._inductor.runtime.triton_compatrj   rh   hashhashlibsha256encode	hexdigestupper)rj   backendkeys      r   triton_hash_with_backendru      sv    @@@@@@GZ\\
,
,GLLNN
,
,C >#**W--..88::@@BBBr   )r
   )	functoolsrn   typingr   cacher^   r	   r   r   r   r&   r,   r1   r4   r=   rA   	lru_cacherC   r_   rh   strru   r3   r   r   <module>r{      sW              D      sCx eCHo     d     $    * t     Y Y Y Y Y t    > T     T4    $ $.D $. $. $. $.N           C# C C C C C Cr   