
    rh                        d dl Z d dlZd dlm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                  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    f/var/www/html/ai-insurance-compliance-backend/venv/lib/python3.12/site-packages/torch/utils/_triton.pyhas_triton_packager      s2    7%%  s   	 	"""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} yy# t
        $ r Y yw xY w)Nr   create_1d_tma_descriptorcreate_2d_tma_descriptorTF)r   r   $triton.tools.experimental_descriptorr   r   r	   r   s     r    has_triton_experimental_host_tmar      s9    !
    s     	,,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_tmar"   .   s9    !    s    	**c                  .    t               xs
 t               S )N)r"   r    r   r   has_triton_tmar%   >   s    02X6V6XXr   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 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   triton.language.extra.cudar'   r(   r	   triton.languager*   )r   r'   r(   r*   s       r   has_triton_tma_devicer-   C   s     JJ##%

002f<MM%%
    B s$   A' '	A32A37A> >	B
	B
c                      t               rYdd l} | j                  j                         r;| j                  j	                         dk\  r| j
                  j                  s	 ddlm} yy# t        $ r Y yw xY w)Nr   r   r)   TF)
r   r   r   r   r   r   r   r,   r*   r	   )r   r*   s     r   has_triton_stable_tma_apir/   c   sf     JJ##%

002f<MM%%B   s   A% %	A10A1c                      t               s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   )get_interface_for_devicedevice_interfacer   c                 P    | j                   j                         j                  dk\  S )N   )Workerget_device_propertiesmajorr2   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)r2   tritons     r   cpu_extra_checkz#has_triton.<locals>.cpu_extra_check   s    0000r   c                      y)NTr$   r8   s    r   _return_truez has_triton.<locals>._return_true   s    r   )r   xpur;   c                  |    j                         D ](  \  } } |       }|j                         s ||      s( y y)NTF)itemsr   )deviceextra_checkr2   r1   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._dynamo.device_interfacer1   r   bool)r9   r?   rA   rH   r1   rG   s       @@r   
has_tritonrK   v   su    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)r   rM   triton.runtime.driverrN   activeget_current_target)rM   rN   targets      r   triton_backendrS      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   rS   hashhashlibsha256encode	hexdigestupper)r   backendkeys      r   triton_hash_with_backendr^      sR    3G\N!GLLN+
,C >>#**W-.88:@@BBr   )	functoolsrW   typingr   cacherJ   r   r   r   r"   r%   r-   	lru_cacher/   rK   rS   strr^   r$   r   r   <module>rd      sT      D   d   $    t   Y Y Y t  > T4  $ .D . .B       C# C Cr   