Ë
    ÌÓ^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Ó5€IØœRŸZ™ZÑ'Ø-¨i¨[¸Ð:‰á-6˜"˜Y˜K qÑ)¸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ÑOùs   ‚ú.é   )r   ÚtupleÚ__version__Úsplitr   )r   r   ÚmajorÚminors       r   Úget_triton_versionr-   .   sT   € ðÛäÑO¨V×-?Ñ-?×-EÑ-EÀcÓ-JÈ2ÈAÐ-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   € ãð 	�
‰
×ÑÓ!ò 	"Ø�J‰J×,Ñ,Ó.°&Ñ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ÒXÔ6VÓ6XÐ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—   € äÔÛð �J‰J×#Ñ#Ô%Ø—
‘
×0Ñ0Ó2°fÒ<Ø—M‘M×%Ò%Ø�Y‰Y×#Ñ#Ô%ð÷ð
 ð øô ò ÙðúðÝ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   € ãð 	�
‰
×ÑÔ!Ø�J‰J×,Ñ,Ó.°'Ò9Ø�J‰J×,Ñ,Ó.°Ò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   € äÔÛð �J‰J×#Ñ#Ô%Ø—
‘
×0Ñ0Ó2°fÒ<Ø—M‘M×%Ò%Ø�Y‰Y×#Ñ#Ô%ðÝ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ò 	ÑˆF�KÙ7¸Ó?ÐØ×,Ñ,Õ.±;Ð?OÕ3PÙð	ð 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)z×Return 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Ó4€KÙØØ€EÜ˜[Ó)ò O‰
ˆˆdÜ�7‰7�>‰>˜$ÕÜ�d˜DÓ!ð O QØ—‘ ˜v Q¤w§~¡~°a·f±f³hÓ'?×'IÑ'IÓ'KÐ&LÐMÔN÷Oð OðOð �8‰8�E‹?Ð÷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˜!˜GŸL™L›NÐ+Ð
,€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Ð )ð	¨S°4©Zó 	ð¨ð ¸$ó ð ‡�ð˜Dò ó ðð ‡�ñ  s¨C x¡ð ¸eÀCÈÀH¹oò ó ðð ‡�ð˜dò ó ðð ‡�ð¨$ò ó ðð* ‡�ð¨tò ó ðð ‡�ðY˜ò Yó ðYð ‡�ð˜tò ó ðð> ‡�ð¨Tò ó ðð €×Ñ�TÓð 4ò ó ðð$ ‡�ð$.�Dò $.ó ð$.ðN ‡�ð
˜ò 
ó ð
ð˜cð  có ð$ ‡�ðC #ò Có ñCr   