
    ^j                     |   d dl Z d dlZd dlZd dlmZ dZdedz  fdZdeddfd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dedefdZe j                  defd       Zy)    N)Anyz(failed to map segment from shared objectreturnc                      t         j                  j                  d      x} r| S 	 ddlm} |j
                  j                  S # t        t        f$ r Y y w xY w)NTRITON_CACHE_DIRr   )knobs)	osenvirongettriton.runtime.cacher   cachedirAttributeErrorImportError)triton_cache_dirr   s     ^/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/utils/_triton.py#_triton_cache_dir_for_error_messager   
   sO    ::>>*<===.{{K( s   A AAexcc                     t               }dt        j                  v rd| d}n
|rd| dnd}t        |  d| d      | )Nr   z (TRITON_CACHE_DIR=)z ( z-. This usually means Triton's cache directoryz is on a filesystem mounted with noexec, so the dynamic loader cannot map generated shared objects. Set TRITON_CACHE_DIR to a directory on an executable filesystem.)r   r   r	   r   )r   	cache_dircache_dir_msgs      r   _raise_triton_cache_load_errorr      s_    35IRZZ'-i[:-6"YKq)B
%<]O L! 	!
     c                  ,    	 dd l } y# t        $ r Y yw xY w)Nr   TF)tritonr   )r   s    r   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>3   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r1   s    r   _device_supports_tmar8   9   sK     	

! 	"JJ,,.&8	"!!!r   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   r8   $triton.tools.experimental_descriptorr:   r;   r<   r   )r:   r;   r<   s      r    has_triton_experimental_host_tmar>   D   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   r8   triton.tools.tensor_descriptorrA   r   r@   s    r   %has_triton_tensor_descriptor_host_tmarC   Z   s9    !    s    	**c                  .    t               xs
 t               S r!   )rC   r>    r   r   has_triton_tmarF   j   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   r1   r2   r3   r4   r5   r6   xputriton.language.extra.cudarH   rI   r   triton.languagerK   )r1   rH   rI   rK   s       r   has_triton_tma_devicerO   o   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)r1   r2   r3   r4   r5   r6   rO   rC   r7   s    r   #has_datacenter_blackwell_tma_devicerS      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/   rJ   TF)r   r1   r2   r3   r4   r5   r6   rL   rN   rK   r   )r1   rK   s     r   has_triton_stable_tma_apirU      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+   rY   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)rY   r   s     r   cpu_extra_checkz#has_triton.<locals>.cpu_extra_check   s    0000r   c                      y)NTrE   r^   s    r   _return_truez has_triton.<locals>._return_true   s    r   )r2   rL   ra   mtiac                  |    j                         D ](  \  } } |       }|j                         s ||      s( y y)NTF)itemsr3   )deviceextra_checkrY   rX   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.configrW   torch._dynamo.device_interfacerX   r   bool)rW   r_   rd   rf   rm   rX   rl   s        @@r   
has_tritonrq      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 # t        t        f$ r"}t        t        |      v rt        |        d }~ww xY w)Nr   )make_backend)driver)triton.compiler.compilerrs   triton.runtime.driverrt   activeget_current_targetr   OSError)_FAILED_TO_MAP_SEGMENT_FROM_SHARED_OBJECTstrr   )rs   rt   targetes       r   triton_backendr~      sV    5,113F##! 4A>*1-s   !0 A!AA!backendc           
         | j                  i       }t        |dd      }|syg }t        |      D ]  \  }}t        j                  j                  |      s&t        |d      5 }|j                  | dt        j                  |j                               j                                 ddd        dj                  |      S # 1 sw Y   xY w)zReturn a cache key fragment for extern libs (e.g. libdevice.10.bc).

    These files affect codegen but are not covered by triton_key() (Python
    sources only) or backend.hash() (ptxas version and arch only).
    extern_libsNr   rb-)parse_optionsgetattrsortedr   pathisfileopenappendhashlibsha256read	hexdigestjoin)r   optsr   partsnamer   fs          r   _extern_libs_keyr      s       $D$t4KE[) O
d77>>$dD! OQvQw~~affh'?'I'I'K&LMNO OO 88E?O Os   !ACC	c                      ddl m}  t               } |         d|j                          }t	        j
                  |j                  d            j                         j                         S )Nr   )
triton_keyr   zutf-8)	%torch._inductor.runtime.triton_compatr   r~   hashr   r   encoder   upper)r   r   keys      r   triton_hash_with_backendr      sR    @G\N!GLLN+
,C >>#**W-.88:@@BBr   ))r   r   )	functoolsr   r   typingr   rz   r{   r   BaseExceptionr   r   rp   r   r(   r"   r-   r8   r>   rC   rF   rO   rS   	lru_cacherU   rq   r~   r   r   rE   r   r   <module>r      s     	  -W )	S4Z 	 $  D   sCx eCHo   d   $  * t   Y Y Y t  > T   T4  $ $.D $. $.N 
 
 
c c $ C# C Cr   