§
    NŠtj1]  ã                   ó,  — d dl mZmZmZ d dlmZmZmZmZ d dl	m
Z
 d dlmZ d dlmZ d dlZd dlmZmZmZmZ d dlmZ d dlZd dlZd dlZd dlZd dlZd dlZd d	lmZ d
efd„Zde de
j!        fd„Z" ej#        ¦   «         dde fd„¦   «         Z$ ej#        ¦   «         de fd„¦   «         Z%de fd„Z& ej#        ¦   «         de fd„¦   «         Z' ej#        d¦  «        d„ ¦   «         Z(de fd„Z) ed¬¦  «         G d„ d¦  «        ¦   «         Z* G d„ de¦  «        Z+dS )é    )ÚBaseBackendÚ	GPUTargetÚLanguage)ÚirÚpassesÚllvmÚnvidia)Úknobs)Ú
PTXASError)Ú	dataclassN)ÚAnyÚDictÚTupleÚOptional)Ú
ModuleType)ÚPathÚtargetc                 óL   — dt           t          t          t          f         fd„}|S )NÚreturnc                 óf   — | j         j        }|j         j        }||k    s
J d¦   «         ‚|dk    rdS dS )Nz%lhs and rhs bitwidth must be the sameé   )é   r   é    )r   r   é   )ÚscalarÚprimitive_bitwidth)Úlhs_typeÚrhs_typeÚlhs_bitwidthÚrhs_bitwidths       ú]/var/www/html/CA-Chatbot/venv/lib/python3.11/site-packages/triton/backends/nvidia/compiler.pyÚcheck_dot_compatibilityz-min_dot_size.<locals>.check_dot_compatibility   sE   € Ø”Ô9ˆØ”Ô9ˆØ˜|Ò+Ð+Ð+Ð-TÑ+Ô+Ð+à˜1ÒÐØ�:à�:ó    )r   Úint)r   r"   s     r!   Úmin_dot_sizer%      s0   € ðµu½SÅ#Ås¸]Ô7Kð ð ð ð ð #Ð"r#   Úarchr   c                 óR   — | dk    rt           j        j        nt           j        j        S )Néd   )r
   r	   Úptxas_blackwellÚptxas)r&   s    r!   Ú	get_ptxasr+   "   s   € Ø+/°3ª;¨;�5Œ<Ô'Ð'½E¼LÔ<NÐNr#   éP   c                 ó¦   — t           j        j        }|�|S t          j        t          | ¦  «        j        dg¦  «                             d¦  «        }|S )Nz	--versionúutf-8)r
   r	   Úmock_ptx_versionÚ
subprocessÚcheck_outputr+   ÚpathÚdecode)r&   Úmock_verÚversions      r!   Úget_ptxas_versionr6   &   sI   € åŒ|Ô,€HØÐØˆÝÔ%¥y°¡¤Ô';¸[Ð&IÑJÔJ×QÒQÐRYÑZÔZ€GØ€Nr#   c                 ó4  — t          | t          ¦  «        sJ ‚t          t          |                      d¦  «        ¦  «        \  }}|dk    r|dk     rd|z   S d|z   dz
  S |dk    rd|z   S |dk    rd	|z   S |d
k    rd}||d
z
  dz  z   |z   S t          d| z   ¦  «        ‚)zK
    Get the highest PTX version supported by the current CUDA driver.
    ú.é   é   r,   r   é   éF   é
   é?   é   éZ   z?Triton only support CUDA 10.0 or higher, but got CUDA version: )Ú
isinstanceÚstrÚmapr$   ÚsplitÚRuntimeError)Úcuda_versionÚmajorÚminorÚbase_ptxs       r!   Úptx_get_versionrJ   /   sÇ   € õ
 �l¥CÑ(Ô(Ð(Ð(Ð(Ý•s˜L×.Ò.¨sÑ3Ô3Ñ4Ô4�L€Eˆ5Ø�‚{€{Ø�1Š9ˆ9Ø˜‘:Ðà˜‘: ‘>Ð!Ø�‚{€{Ø�E‰zÐØ�‚{€{Ø�E‰zÐà�‚{€{ØˆØ˜5 2™:¨Ñ+Ñ+¨eÑ3Ð3å
ÐXÐ[gÑgÑ
hÔ
hÐhr#   c                 ó^   — | j         }|€#t          |¦  «        j        }t          |¦  «        }|S ©N)Úptx_versionr+   r5   rJ   )Úoptionsr&   rM   rF   s       r!   Úget_ptx_version_from_optionsrO   G   s1   € ØÔ%€KØÐÝ  ‘”Ô.ˆÝ% lÑ3Ô3ˆØÐr#   c                 óP   — t          | |¦  «        }t          d|¦  «        }d|› �}|S )Nr@   z+ptx)rO   Úmin)rN   r&   rM   Úllvm_ptx_versionÚfeaturess        r!   Úget_featuresrT   O   s6   € å.¨w¸Ñ=Ô=€Kõ ˜2˜{Ñ+Ô+ÐØ(Ð&Ð(Ð(€HØ€Or#   c                 óÆ   — t          | d¦  «        5 }t          j        |                     ¦   «         ¦  «                             ¦   «         cd d d ¦  «         S # 1 swxY w Y   d S )NÚrb)ÚopenÚhashlibÚsha256ÚreadÚ	hexdigest)r2   Úfs     r!   Ú	file_hashr]   ]   s”   € å	ˆd�DÑ	Ô	ð 4˜QÝŒ~˜aŸfšf™hœhÑ'Ô'×1Ò1Ñ3Ô3ð4ð 4ð 4ð 4ñ 4ô 4ð 4ð 4ð 4ð 4ð 4ð 4øøøð 4ð 4ð 4ð 4ð 4ð 4s   ‘8AÁAÁAÚ
capabilityc                 ó$   — | dk    rdnd}d| › |› �S )Nr@   ÚaÚ Úsm_© )r^   Úsuffixs     r!   Úsm_arch_from_capabilityre   c   s(   € à "Ò$Ð$ˆSˆS¨"€FØ%�Ð%˜VÐ%Ð%Ð%r#   T)Úfrozenc                   óÐ  — e Zd ZU dZeed<   dZeed<   dZeed<   dZeed<   d	Z	e
e         ed
<   d	Zeed<   ej        j        Ze
e         ed<   d	Ze
e         ed<   dZeed<   dZeed<   dZeed<   dZeed<   dZee         ed<   dZee         ed<   dZeed<   dZee         ed<   d	Zeed<   d	Zeed<   dZeed<   dZ eed <   dZ!eed!<   d	Z"eed"<   d#Z#eed$<   d%„ Z$d&„ Z%e&d'„ ¦   «         Z'd	S )(ÚCUDAOptionsé   Ú	num_warpsr   Únum_ctasé   Ú
num_stagesr   Ú	warp_sizeNÚmaxnregrM   Úptx_optionsÚir_overrideTÚenable_fp_fusionÚenable_reflect_ftzFÚlaunch_cooperative_gridÚ
launch_pdl)Úfp8e5Úfp8e4b15Úsupported_fp8_dtypesrc   Ú!deprecated_fp8_dot_operand_dtypesÚtf32Údefault_dot_input_precision)rz   Útf32x3ÚieeeÚbf16x3Úbf16x6Úallowed_dot_input_precisionsÚmax_num_imprecise_acc_defaultÚextern_libsÚdebugÚcudaÚbackend_nameÚsanitize_overflowr&   ra   Úinstrumentation_modec                 ó¼  — t          t          ¦  «        j        dz  }| j        €i nt	          | j        ¦  «        }|                     dd ¦  «        s&t          j        j        pt          |dz  ¦  «        |d<   t                               | dt          |                     ¦   «         ¦  «        ¦  «         | j        dk    r| j        | j        dz
  z  dk    s
J d¦   «         ‚d S )NÚlibÚ	libdevicezlibdevice.10.bcr‚   r   r   znum_warps must be a power of 2)r   Ú__file__Úparentr‚   ÚdictÚgetr
   r	   Úlibdevice_pathrB   ÚobjectÚ__setattr__ÚtupleÚitemsrj   )ÚselfÚdefault_libdirr‚   s      r!   Ú__post_init__zCUDAOptions.__post_init__…   sÔ   € Ý�h™œÔ.°Ñ6ˆØ Ô,Ð4�b�b½$¸tÔ?OÑ:PÔ:PˆØ�Š˜{¨DÑ1Ô1ð 	nÝ',¤|Ô'BÐ'mÅcÈ.Ð[lÑJlÑFmÔFmˆK˜Ñ$å×Ò˜4 µ°k×6GÒ6GÑ6IÔ6IÑ0JÔ0JÑKÔKÐKØŒ~ Ò!Ð! t¤~¸¼È!Ñ9KÑ'LÐQRÒ&RÐ&RÐ&RØ/ñ 'SÔ&RÐRÐ&RÐ&Rr#   c                 óv  — t          | j        ¦  «        }t          d„ t          |d         ¦  «        D ¦   «         ¦  «        |d<   d                     d„ t          |                     ¦   «         ¦  «        D ¦   «         ¦  «        }t          j        |                     d¦  «        ¦  «         	                    ¦   «         S )Nc              3   ó>   K  — | ]\  }}|t          |¦  «        fV — Œd S rL   )r]   )Ú.0ÚkÚvs      r!   ú	<genexpr>z#CUDAOptions.hash.<locals>.<genexpr>‘   s1   è è € Ð(hÐ(h¹t¸qÀ!¨!­Y°q©\¬\Ð):Ð(hÐ(hÐ(hÐ(hÐ(hÐ(hr#   r‚   Ú_c                 ó"   — g | ]\  }}|› d |› �‘ŒS )ú-rc   )r™   ÚnameÚvals      r!   ú
<listcomp>z$CUDAOptions.hash.<locals>.<listcomp>’   s&   € ÐSÐSÐS©I¨D°#˜4˜˜ #˜˜ÐSÐSÐSr#   r.   )
r�   Ú__dict__r’   ÚsortedÚjoinr“   rX   rY   Úencoder[   )r”   Ú	hash_dictÚkeys      r!   ÚhashzCUDAOptions.hash�   sš   € Ý˜œÑ'Ô'ˆ	Ý#(Ð(hÐ(hÅvÈiÐXeÔNfÑGgÔGgÐ(hÑ(hÔ(hÑ#hÔ#hˆ	�-Ñ Ø�hŠhÐSÐS½À	ÇÂÑ@QÔ@QÑ9RÔ9RÐSÑSÔSÑTÔTˆÝŒ~˜cŸjšj¨Ñ1Ô1Ñ2Ô2×<Ò<Ñ>Ô>Ð>r#   c                 ó   — d| j         v S )NÚiisan)r‡   )r”   s    r!   Úenable_iisanzCUDAOptions.enable_iisan•   s   € à˜$Ô3Ð3Ð3r#   )(Ú__name__Ú
__module__Ú__qualname__rj   r$   Ú__annotations__rk   rm   rn   ro   r   rM   r
   r	   Úptxas_optionsrp   rB   rq   rr   Úboolrs   rt   ru   rx   r   ry   r{   r€   r�   r‚   r�   rƒ   r…   r†   r&   r‡   r–   r©   Úpropertyr¬   rc   r#   r!   rh   rh   i   sè  € € € € € € à€IˆsÐÐÑØ€HˆcÐÐÑØ€J�ÐÐÑØ€IˆsÐÐÑð "€GˆX�cŒ]Ð!Ð!Ñ!Ø€K�ÐÐÑØ!&¤Ô!;€K�˜#”Ð;Ð;Ñ;Ø!%€K�˜#”Ð%Ð%Ñ%Ø!Ð�dÐ!Ð!Ñ!Ø#Ð˜Ð#Ð#Ñ#Ø$)Ð˜TÐ)Ð)Ñ)Ø€J�ÐÐÑØ'<Ð˜% œ*Ð<Ð<Ñ<Ø46Ð% u¨S¤zÐ6Ð6Ñ6Ø'-Ð Ð-Ð-Ñ-Ø/]Ð  %¨¤*Ð]Ð]Ñ]Ø*.Ð! 4Ð.Ð.Ñ.Ø€K�ÐÐÑØ€Eˆ4ÐÐÑØ€L�#ÐÐÑØ"Ð�tÐ"Ð"Ñ"Ø€Dˆ#ÐÐÑØ "Ð˜#Ð"Ð"Ñ"ð0ð 0ð 0ð?ð ?ð ?ð ð4ð 4ñ „Xð4ð 4ð 4r#   rh   c                   ó  ‡ — e Zd ZdZedefd„¦   «         Zd„ Zdefd„Z	deddfˆ fd„Z
defd„Zd	„ Zd
„ Zdeeef         fd„Zd„ Zed„ ¦   «         Zed„ ¦   «         Zd„ Zd„ Zd„ Zd„ Zd„ Z ej        ¦   «         d„ ¦   «         Zˆ xZS )ÚCUDABackendNr   c                 ó   — | j         dk    S )Nr„   )Úbackend)r   s    r!   Úsupports_targetzCUDABackend.supports_target�   s   € àŒ~ Ò'Ð'r#   c                 óœ   — d}t          j        ||¦  «        }|st          d|› �¦  «        ‚t          |                     d¦  «        ¦  «        S )Nz	^sm(\d+)$z(TRITON_OVERRIDE_ARCH must have the form r   )ÚreÚ	fullmatchÚ
ValueErrorr$   Úgroup)r”   r&   ÚpatternÚmatchs       r!   Ú_parse_archzCUDABackend._parse_arch¡   sP   € ØˆÝ”˜W dÑ+Ô+ˆØð 	SÝÐQÈÐQÐQÑRÔRÐRÝ�5—;’;˜q‘>”>Ñ"Ô"Ð"r#   r   c                 ó@   — |                       |j        ¦  «        }d|› �S )Núcuda:)rÀ   r&   )r”   rN   r^   s      r!   Úget_target_namezCUDABackend.get_target_name¨   s%   € Ø×%Ò% g¤lÑ3Ô3ˆ
Ø#�zÐ#Ð#Ð#r#   c                 óX   •— t          ¦   «                              |¦  «         d| _        d S )NÚcubin)ÚsuperÚ__init__Ú
binary_ext)r”   r   Ú	__class__s     €r!   rÇ   zCUDABackend.__init__¬   s&   ø€ Ý‰Œ×Ò˜Ñ Ô Ð Ø!ˆŒˆˆr#   c                 óð  ‡— t          ˆfd„dD ¦   «         ¦  «        r
d‰d<   d‰d<   dt          j        j        pd| j        j        › �i}|                     ˆfd	„t          j         	                    ¦   «         D ¦   «         ¦  «         t          |                      |d         ¦  «        ¦  «        }|                     d
d¦  «        dk    r|dk     rt          d|› d�¦  «        ‚d|vrSt          t          j        ¦  «        }|dk    r|                     d¦  «         t#          t%          |¦  «        ¦  «        |d<   d|vr|dk    rd|d<   d|vrt          j        j        |d<   |dk    rdnd|d<   t          di |¤ŽS )Nc              3   óH   •K  — | ]}|‰                      d d¦  «        v V — ŒdS )r‡   ra   N)rŽ   )r™   ÚmodeÚoptss     €r!   rœ   z,CUDABackend.parse_options.<locals>.<genexpr>²   s7   øè è € Ð\Ð\Àˆt�t—x’xÐ 6¸Ñ;Ô;Ð;Ð\Ð\Ð\Ð\Ð\Ð\r#   )Úconsanr«   Trƒ   Fr†   r&   Úsmc                 ó:   •— i | ]}|‰v ¯‰|         ®|‰|         “ŒS rL   rc   )r™   rš   rÍ   s     €r!   ú
<dictcomp>z-CUDABackend.parse_options.<locals>.<dictcomp>·   s7   ø€ ÐuÐuÐu AÐTUÐY]ÐT]ÐT]ÐaeÐfgÔahÐat�Q˜˜QœÐatÐatÐatr#   rk   r   r@   zBnum_ctas > 1 requires NVIDIA SM90+ (Hopper). Current target is sm_zM. This configuration will fail. Please set num_ctas=1 or target an SM90+ GPU.rx   éY   Úfp8e4nvry   )rw   rr   i   @r   r�   rc   )Úanyr
   ÚruntimeÚoverride_archr   r&   Úupdaterh   Ú__dataclass_fields__Úkeysr$   rÀ   rŽ   r¼   Úsetrx   Úaddr’   r¤   ÚlanguageÚdefault_fp_fusion)r”   rÍ   Úargsr^   rx   s    `   r!   Úparse_optionszCUDABackend.parse_options°   sÇ  ø€ åÐ\Ð\Ð\Ð\ÐH[Ð\Ñ\Ô\Ñ\Ô\ð 	.Ø ˆD�‰MØ(-ˆDÐ$Ñ%à�œÔ3ÐNÐ7N¸D¼KÔ<LÐ7NÐ7NÐOˆØ�ŠÐuÐuÐuÐu­Ô)I×)NÒ)NÑ)PÔ)PÐuÑuÔuÑvÔvÐvÝ˜×)Ò)¨$¨v¬,Ñ7Ô7Ñ8Ô8ˆ
à�8Š8�J Ñ"Ô" QÒ&Ð&¨:¸ª?¨?Ýð OØ6@ðOð Oð Oñ Qô Qð Qð "¨Ð-Ð-Ý#&¥{Ô'GÑ#HÔ#HÐ Ø˜RÒÐØ$×(Ò(¨Ñ3Ô3Ð3Ý+0µÐ8LÑ1MÔ1MÑ+NÔ+NˆDÐ'Ñ(à.°dÐ:Ð:Ø˜RÒÐØ<J�Ð8Ñ9à TÐ)Ð)Ý',¤~Ô'GˆDÐ#Ñ$à9CÀrÒ9IÐ9I°°ÈqˆÐ,Ñ-åÐ"Ð"˜TÐ"Ð"Ð"r#   c                 ó*   — |j         |j        |j        fS rL   )rj   rk   Úshared)r”   Úmetadatas     r!   Úpack_metadatazCUDABackend.pack_metadataÐ   s   € àÔØÔØŒOð
ð 	
r#   c                 ó¾   — dd l mc mc m} t	          |                      |j        ¦  «        ¦  «        }|dk    r|j        n|j        t          | j
        ¦  «        dœ}|S )Nr   r,   )Úconvert_custom_typesr%   )Útriton.language.extra.cudarÜ   Úextrar„   r$   rÀ   r&   Úconvert_custom_float8_sm80Úconvert_custom_float8_sm70r%   r   )r”   rN   r„   r^   Úcodegen_fnss        r!   Úget_codegen_implementationz&CUDABackend.get_codegen_implementation×   s~   € Ø1Ð1Ð1Ð1Ð1Ð1Ð1Ð1Ð1Ð1Ð1Ð1Ý˜×)Ò)¨'¬,Ñ7Ô7Ñ8Ô8ˆ
ð 0:¸RÒ/?Ð/?ˆDÔ+Ð+ÀTÔEdÝ˜œÑ%Ô%ð
ð 
ˆð
 Ðr#   c                 ó   — ddl m} d|iS )Nr   )rŠ   ztriton.language.extra.libdevice)ræ   rŠ   )r”   rŠ   s     r!   Úget_module_mapzCUDABackend.get_module_mapá   s   € Ø8Ð8Ð8Ð8Ð8Ð8Ø1°9Ð=Ð=r#   c                 óˆ   — t          j        |¦  «         t          j        r!t          j                             |¦  «         d S d S rL   )r	   Úload_dialectsrµ   Úinstrumentation)r”   Úctxs     r!   rï   zCUDABackend.load_dialectså   sE   € ÝÔ˜SÑ!Ô!Ð!ÝÔ&ð 	;ÝÔ'×5Ò5°cÑ:Ô:Ð:Ð:Ð:ð	;ð 	;r#   c                 óÌ  — t          j        | j        ¦  «        }|                     ¦   «          t          j                             |¦  «         t          j                             |¦  «         |dz  dk     rt          j         	                    |¦  «         t          j         
                    |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         |                     | d¦  «         | S )Nr=   é	   Ú	make_ttir)r   Úpass_managerÚcontextÚenable_debugr   ÚcommonÚadd_inlinerÚttirÚadd_rewrite_tensor_pointerÚ(add_rewrite_tensor_descriptor_to_pointerÚadd_canonicalizerÚadd_combineÚadd_reorder_broadcastÚadd_cseÚadd_symbol_dceÚadd_loop_unrollÚrun)Úmodrâ   Úoptr^   Úpms        r!   rô   zCUDABackend.make_ttirê   s  € åŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐÝŒ×!Ò! "Ñ%Ô%Ð%ÝŒ×.Ò.¨rÑ2Ô2Ð2Ø˜Ñ˜aÒÐÝŒK×@Ò@ÀÑDÔDÐDÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò Ñ#Ô#Ð#ÝŒ×)Ò)¨"Ñ-Ô-Ð-ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒ×$Ò$ RÑ(Ô(Ð(ÝŒ×#Ò# BÑ'Ô'Ð'Ø
�Šˆs�KÑ Ô Ð Øˆ
r#   c                 óf  — |j         �E|                      dt          j        | j        ¦  «                             |j         ¦  «        ¦  «         t          j        | j        ¦  «        }|                     ¦   «         }|dz  dk    }t          j	         
                    |d|› �|j        d|j        ¦  «         t          j                             |¦  «         t          j                             ||¦  «         t           j        j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             ||dk    ¦  «         t           j        j                             |¦  «         t          j	                             |¦  «         |dz  dv �r2t          j                             |¦  «         t          j                             |¦  «         t          j	                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t           j        j                             ||j         |¦  «         t          j         !                    ||j         ¦  «         t          j         "                    |¦  «         t          j         #                    ||j         |¦  «         �nû|dz  dk    �rÒt          j                             |¦  «         t          j                             |¦  «         t          j	                             |¦  «         t          j         $                    |¦  «         t          j         %                    |d¦  «         t           j        j         &                    |¦  «         t          j         !                    ||j         ¦  «         t          j         "                    |¦  «         t          j         '                    ||j         ¦  «         t          j         #                    ||j         |¦  «         t          j         (                    |¦  «         t          j                             |¦  «         t          j         %                    |d	¦  «         t           j        j         )                    |¦  «         nt          j	                             |¦  «         t          j                             |¦  «         t          j	                             |¦  «         t          j         *                    |¦  «         t          j                             ||dk    ¦  «         t          j         +                    |¦  «         t           j        j         ,                    |¦  «         |dz  d
k    r$t           j        j         -                    |¦  «         t          j                             |¦  «         t           j        j         .                    |¦  «         t          j         /                    |¦  «         t          j         0                    |¦  «         t          j	                             |¦  «         t          j         1                    |¦  «         t           j        j         2                    ||¦  «         t           j        j         3                    |¦  «         t          j         4                    |¦  «         t          j         5                    |¦  «         t          j                             |¦  «         | 6                    | d¦  «         |  7                    ¦   «         |d<   | S )Nzttg.maxnregr=   r   rÂ   r   r,   )r   ró   FTró   Ú
make_ttgirÚtensordesc_meta)8ro   Úset_attrr   Úbuilderrö   Úget_int32_attrrõ   r÷   r   rú   Úadd_convert_to_ttgpuirrj   rk   ÚttgpuirÚadd_coalesceÚadd_f32_dot_tcr	   Ú	ttnvgpuirÚadd_plan_ctaÚadd_remove_layout_conversionsÚadd_optimize_thread_localityÚadd_accelerate_matmulÚadd_optimize_dot_operandsÚ add_optimize_descriptor_encodingÚadd_loop_aware_cseÚadd_fuse_nested_loopsrø   rý   Úadd_triton_licmÚ add_combine_tensor_select_and_ifÚhopperÚadd_hopper_warpspecrm   Úadd_assign_latenciesÚadd_schedule_loopsÚadd_pipelineÚadd_optimize_accumulator_initÚadd_hoist_tmem_allocÚadd_promote_lhs_to_tmemÚadd_warp_specializeÚadd_optimize_partition_warpsÚadd_remove_tmem_tokensÚadd_prefetchÚadd_coalesce_async_copyÚadd_optimize_tmem_layoutsÚadd_tma_loweringÚadd_interleave_tmemÚadd_reduce_data_duplicationÚadd_reorder_instructionsr  Úadd_fence_insertionÚadd_lower_mmaÚadd_sccpr   r  Úget_tensordesc_metadata)r  râ   r  r^   r  Údump_enabledÚemuTF32s          r!   r  zCUDABackend.make_ttgirû   sÉ  € ð Œ;Ð"Ø�LŠL˜­¬
°3´;Ñ(?Ô(?×(NÒ(NÈsÌ{Ñ([Ô([Ñ\Ô\Ð\åŒ_˜Sœ[Ñ)Ô)ˆØ—’Ñ(Ô(ˆØ Ñ# qÒ(ˆÝŒ×*Ò*¨2Ð/C°zÐ/CÐ/CÀSÄ]ÐTVÐX[ÔXdÑeÔeÐeåŒ×#Ò# BÑ'Ô'Ð'ÝŒ×%Ò% b¨'Ñ2Ô2Ð2åŒÔ×,Ò,¨RÑ0Ô0Ð0ÝŒ×4Ò4°RÑ8Ô8Ð8ÝŒ×3Ò3°BÑ7Ô7Ð7ÝŒ×,Ò,¨RÑ0Ô0Ð0ÝŒ×4Ò4°RÑ8Ô8Ð8ÝŒ×0Ò0°°ZÀ2Ò5EÑFÔFÐFÝŒÔ×@Ò@ÀÑDÔDÐDÝŒ×&Ò& rÑ*Ô*Ð*Ø˜Ñ˜vÐ%Ñ%ÝŒN×0Ò0°Ñ4Ô4Ð4ÝŒM×+Ò+¨BÑ/Ô/Ð/ÝŒK×'Ò'¨Ñ+Ô+Ð+ÝŒM×+Ò+¨BÑ/Ô/Ð/ÝŒN×;Ò;¸BÑ?Ô?Ð?ÝŒMÔ ×4Ò4°R¸¼ÈÑVÔVÐVÝŒN×/Ò/°°C´NÑCÔCÐCÝŒN×-Ò-¨bÑ1Ô1Ð1ÝŒN×'Ò'¨¨C¬N¸LÑIÔIÐIÑIØ˜2Ñ Ò#Ñ#ÝŒN×0Ò0°Ñ4Ô4Ð4ÝŒM×+Ò+¨BÑ/Ô/Ð/ÝŒK×'Ò'¨Ñ+Ô+Ð+ÝŒN×8Ò8¸Ñ<Ô<Ð<ÝŒN×/Ò/°°EÑ:Ô:Ð:ÝŒMÔ#×;Ò;¸BÑ?Ô?Ð?ÝŒN×/Ò/°°C´NÑCÔCÐCÝŒN×-Ò-¨bÑ1Ô1Ð1ÝŒN×.Ò.¨r°3´>ÑBÔBÐBÝŒN×'Ò'¨¨C¬N¸LÑIÔIÐIÝŒN×7Ò7¸Ñ;Ô;Ð;ÝŒN×;Ò;¸BÑ?Ô?Ð?åŒN×/Ò/°°DÑ9Ô9Ð9ÝŒMÔ#×:Ò:¸2Ñ>Ô>Ð>Ð>åŒK×'Ò'¨Ñ+Ô+Ð+ÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×&Ò& rÑ*Ô*Ð*ÝŒ×#Ò# BÑ'Ô'Ð'ÝŒ×0Ò0°°ZÀ2Ò5EÑFÔFÐFÝŒ×.Ò.¨rÑ2Ô2Ð2ÝŒÔ×9Ò9¸"Ñ=Ô=Ð=Ø˜Ñ˜qÒ Ð ÝŒMÔ#×4Ò4°RÑ8Ô8Ð8ÝŒ×4Ò4°RÑ8Ô8Ð8ÝŒÔ×3Ò3°BÑ7Ô7Ð7ÝŒ×2Ò2°2Ñ6Ô6Ð6ÝŒ×/Ò/°Ñ3Ô3Ð3ÝŒ×&Ò& rÑ*Ô*Ð*ÝŒ×$Ò$ RÑ(Ô(Ð(ÝŒÔ×3Ò3°B¸
ÑCÔCÐCÝŒÔ×-Ò-¨bÑ1Ô1Ð1ÝŒ×Ò˜rÑ"Ô"Ð"ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒ×'Ò'¨Ñ+Ô+Ð+à
�Šˆs�LÑ!Ô!Ð!Ø&)×&AÒ&AÑ&CÔ&CˆÐ"Ñ#Øˆ
r#   c                 óö  — |}t          j        |j        ¦  «        }|                     ¦   «          t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j        j
                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         |                     |d¦  «         |                     ¦   «         |d<   |S )NÚgluon_to_ttgirr	  )r   rõ   rö   r÷   r   Úgluonrù   Úadd_infer_coalesced_encodingsÚadd_resolve_auto_encodingsr	   r  r*  rý   rø   r0  rú   r  r  r  r  r1  )r”   Úsrcrâ   rN   r^   r  r  s          r!   r5  zCUDABackend.gluon_to_ttgirE  s"  € ØˆÝŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐåŒ× Ò  Ñ$Ô$Ð$ÝŒ×2Ò2°2Ñ6Ô6Ð6ÝŒ×/Ò/°Ñ3Ô3Ð3ÝŒÔ×0Ò0°Ñ4Ô4Ð4ÝŒ×&Ò& rÑ*Ô*Ð*ÝŒ×Ò˜rÑ"Ô"Ð"ÝŒ×&Ò& rÑ*Ô*Ð*ÝŒ×&Ò& rÑ*Ô*Ð*ÝŒ×7Ò7¸Ñ;Ô;Ð;à
�ŠˆsÐ$Ñ%Ô%Ð%Ø&)×&AÒ&AÑ&CÔ&CˆÐ"Ñ#Øˆ
r#   c                 ó®  — t          || j        j        ¦  «        }|}t          j        |j        ¦  «        }|                     ¦   «          t          j         	                    |¦  «         t          j         
                    |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j        j                             |||¦  «         t          j        j                             |¦  «         t          j        j                             |¦  «         d|j        v rt          j                             |¦  «         t          j                             |¦  «         t          j        j                             ||¦  «         t0          j        r&t0          j                             d||j        ¦  «         t          j        j                             |||¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j        j                             |¦  «         t          j        j                              |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j         !                    |¦  «         t          j         "                    |¦  «         tF          j$        j%        s0tF          j$        j&        st          j'         (                    |¦  «         t0          j        r&t0          j                             d||j        ¦  «         | )                    |d¦  «         tF          j$        j&        rÕtF          j$        j%        sbt          j        |j        ¦  «        }|                     ¦   «          t          j'         (                    |¦  «         | )                    |d¦  «         t          j        |j        ¦  «        }|                     ¦   «          t          j'         *                    |¦  «         | )                    |d¦  «         tW          j,        ¦   «          tW          j        ¦   «         }tF          j$        j-        rt]          d¦  «        ‚tW          j/        ||¦  «        }	ta          |¦  «        }
tc          || j        j        ¦  «        }d}t          j2        ¦   «          tW          j3        |	||
|¦  «         |j4        rt          j5        |	¦  «         |j6        r:t          j7        |	¦  «        r&d	„ |j6        D ¦   «         }tW          j8        |	|¦  «         tW          j9        |	tV          j:        ¦  «         | ;                    d
¦  «        }|�||d<   | ;                    d¦  «        |d<   | ;                    d¦  «        |d<   | ;                    d¦  «        |d<   | ;                    d¦  «        |d<   | ;                    d¦  «        pd|d<   | ;                    d¦  «        pd|d<   ty          |	¦  «        }~	~|S )NrÎ   Úttgpuir_to_llvmirÚllvmir_to_llvmÚ	make_llirzmake_llir.disable_line_infoz,make_llir.dump_ir_extract_di_local_variableszYAddress Sanitizer Error: Address sanitizer is currently only supported on the AMD backendúnvptx64-nvidia-cudac                 ó   — g | ]\  }}|‘ŒS rc   rc   )r™   r    r2   s      r!   r¢   z)CUDABackend.make_llir.<locals>.<listcomp>£  s   € ÐBÐBÐB™l˜t T�TÐBÐBÐBr#   zttg.total-num-warpsrj   z
ttg.sharedrá   zttg.tensor_memory_sizeÚ	tmem_sizezttg.global_scratch_memory_sizeÚglobal_scratch_sizez#ttg.global_scratch_memory_alignmentÚglobal_scratch_alignzttg.profile_scratch_memory_sizer   Úprofile_scratch_sizez$ttg.profile_scratch_memory_alignmentr   Úprofile_scratch_align)=rO   r   r&   r   rõ   rö   r÷   r   r  r  Úadd_allocate_warp_groupsÚconvertÚadd_scf_to_cfr6  rù   r	   Úadd_allocate_shared_memory_nvr  Úadd_allocate_tensor_memoryÚadd_check_matmul_two_ctar‡   Úadd_concurrency_sanitizerÚ"add_allocate_global_scratch_memoryÚadd_proxy_fence_insertionrµ   rð   ÚpatchÚadd_to_llvmirrø   rý   r   Úadd_nvgpu_to_llvmÚadd_warp_specialize_to_llvmr  Úadd_nvvm_to_llvmr
   ÚcompilationÚdisable_line_infoÚ"dump_ir_extract_di_local_variablesÚllvmirÚadd_di_scoper  Úadd_di_local_variabler   Úinit_targetsÚenable_asanrE   Ú	to_modulere   rT   Úset_short_ptrÚattach_datalayoutrs   Úset_nvvm_reflect_ftzr‚   Úhas_extern_depsÚlink_extern_libsÚoptimize_moduleÚOPTIMIZE_O3Úget_int_attrrB   )r”   r9  râ   rN   r^   rM   r  r  rö   Úllvm_modÚprocrS   ÚtripleÚpathsÚtotal_num_warpsÚrets                   r!   r=  zCUDABackend.make_llirX  sT  € Ý2°7¸D¼KÔ<LÑMÔMˆàˆåŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐåŒ×7Ò7¸Ñ;Ô;Ð;ÝŒ×/Ò/°Ñ3Ô3Ð3ÝŒ×$Ò$ RÑ(Ô(Ð(ÝŒ× Ò  Ñ$Ô$Ð$ÝŒÔ×;Ò;¸BÀ
ÈKÑXÔXÐXÝŒÔ×:Ò:¸2Ñ>Ô>Ð>ÝŒÔ×8Ò8¸Ñ<Ô<Ð<Ø�wÔ3Ð3Ð3åŒN×4Ò4°RÑ8Ô8Ð8ÝŒ×9Ò9¸"Ñ=Ô=Ð=ÝŒÔ×9Ò9¸"¸jÑIÔIÐIåÔ&ð 	TÝÔ'×-Ò-Ð.AÀ2ÀsÄ{ÑSÔSÐSÝŒÔ×+Ò+¨B°
¸KÑHÔHÐHÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒÔ×1Ò1°"Ñ5Ô5Ð5ÝŒÔ×;Ò;¸BÑ?Ô?Ð?ÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒ×$Ò$ RÑ(Ô(Ð(ÝŒ×'Ò'¨Ñ+Ô+Ð+åÔ Ô2ð 	+½5Ô;LÔ;oð 	+ÝŒM×&Ò& rÑ*Ô*Ð*åÔ&ð 	QÝÔ'×-Ò-Ð.>ÀÀCÄKÑPÔPÐPà
�Šˆs�KÑ Ô Ð åÔÔ?ð 	HåÔ$Ô6ð ;Ý”_ S¤[Ñ1Ô1�Ø—’Ñ!Ô!Ð!Ý”×*Ò*¨2Ñ.Ô.Ð.Ø—’�sÐ9Ñ:Ô:Ð:õ ” ¤Ñ-Ô-ˆBØ�OŠOÑÔÐÝŒM×/Ò/°Ñ3Ô3Ð3Ø�FŠF�3ÐFÑGÔGÐGõ 	ÔÑÔÐÝ”,‘.”.ˆÝÔÔ(ð 	mÝØkñmô mð må”> # wÑ/Ô/ˆÝ& zÑ2Ô2ˆÝ ¨¬Ô)9Ñ:Ô:ˆØ&ˆÝÔÑÔÐÝÔ˜x¨°°xÑ@Ô@Ð@ØÔ%ð 	2ÝÔ'¨Ñ1Ô1Ð1àÔð 	3¥6Ô#9¸(Ñ#CÔ#Cð 	3ØBÐB¨gÔ.AÐBÑBÔBˆEÝÔ! (¨EÑ2Ô2Ð2åÔ˜X¥tÔ'7Ñ8Ô8Ð8ð ×*Ò*Ð+@ÑAÔAˆØÐ&Ø$3ˆH�[Ñ!Ø ×-Ò-¨lÑ;Ô;ˆ�ÑØ #× 0Ò 0Ð1IÑ JÔ Jˆ�ÑØ*-×*:Ò*:Ð;[Ñ*\Ô*\ˆÐ&Ñ'Ø+.×+;Ò+;Ð<aÑ+bÔ+bˆÐ'Ñ(Ø+.×+;Ò+;Ð<]Ñ+^Ô+^Ð+cÐbcˆÐ'Ñ(Ø,/×,<Ò,<Ð=cÑ,dÔ,dÐ,iÐhiˆÐ(Ñ)Ý�(‰mŒmˆØØØˆ
r#   c           	      ó˜  — t          || j        j        ¦  «        }d}t          |¦  «        }t	          || j        j        ¦  «        }dg}	t          j        |||||	|j        d¦  «        }
t          j	        d|
¦  «        }t          |¦  «        dk    sJ ‚|d         |d<   |dz  › d	|dz  › �}t          j        d
d|› �|
t          j        ¬¦  «        }
t          j        dd|› �|
t          j        ¬¦  «        }
t          j        j        st          j        dd|
¦  «        }
t          j        j        rt%          d¦  «         t%          |
¦  «         |
S )Nr>  znvptx-mad-wide-optFz(.visible .entry ([a-zA-Z_][a-zA-Z0-9_]*)r   r   r    r=   r8   z\.version \d+\.\d+z	.version )Úflagsz\.target sm_\d+z.target sm_z,\s*debug|debug,\s*ra   z // -----// NVPTX Dump //----- //)rO   r   r&   re   rT   r   Útranslate_to_asmrr   rº   ÚfindallÚlenÚsubÚ	MULTILINEr
   rS  rU  r	   Ú
dump_nvptxÚprint)r”   r9  râ   r  r^   rM   rf  re  rS   rk  ri  Únamess               r!   Úmake_ptxzCUDABackend.make_ptx¸  sN  € Ý2°3¸¼Ô8HÑIÔIˆà&ˆÝ& zÑ2Ô2ˆÝ  T¤[Ô%5Ñ6Ô6ˆØ%Ð&ˆÝÔ# C¨°°xÀÈÔH\Ð^cÑdÔdˆå”
ÐFÈÑLÔLˆÝ�5‰zŒz˜QŠˆˆˆØ  œ8ˆ�Ñà$ b™Ð;Ð;¨;°r©>Ð;Ð;ˆÝŒfÐ*Ð,E¸Ð,EÐ,EÀsÕRTÔR^Ð_Ñ_Ô_ˆÝŒfÐ'Ð)C°zÐ)CÐ)CÀSÕPRÔP\Ð]Ñ]Ô]ˆÝÔ ÔCð 	:õ ”&Ð/°°SÑ9Ô9ˆCÝŒ<Ô"ð 	ÝÐ4Ñ5Ô5Ð5Ý�#‰JŒJˆJØˆ
r#   c           
      óœ  — t          | j        j        ¦  «        j        }t	          j        ddd¬¦  «        5 }t	          j        ddd¬¦  «        5 }|                     |¦  «         |                     ¦   «          |j        dz   }g }	t          j
        j        r|	dd	gz  }	nt          j        j        r|	d
gz  }	n|	dgz  }	|j        rg ndg}
t          |¦  «        }t          j        j        rddgng }|j        r|j                             d¦  «        ng }dg}|g|	¢|
¢d‘|¢|¢|¢d|› �‘|j        ‘d‘|‘}	 t%          j        |dd|¬¦  «         t          j        j        rMt+          |j        ¦  «        5 }t-          |                     ¦   «         ¦  «         d d d ¦  «         n# 1 swxY w Y   t0          j                             |j        ¦  «        rt1          j        |j        ¦  «         t0          j                             |j        ¦  «        rt1          j        |j        ¦  «         �n
# t$          j        $ rø}t+          |j        ¦  «        5 }|                     ¦   «         }d d d ¦  «         n# 1 swxY w Y   t0          j                             |j        ¦  «        rt1          j        |j        ¦  «         |j        dk    rd}n%|j        dt:          j        z   k    rd}n
d|j        › �}|› d|› dd                     |¦  «        › d�}t-          d|› d|› d�¦  «         tA          |¦  «        ‚d }~ww xY wt+          |d ¦  «        5 }|                     ¦   «         }d d d ¦  «         n# 1 swxY w Y   t0          j                             |¦  «        rt1          j        |¦  «         d d d ¦  «         n# 1 swxY w Y   d d d ¦  «         n# 1 swxY w Y   |S )!NFÚwz.ptx)ÚdeleterÌ   rd   Úrz.logz.oz	-lineinfoz-suppress-debug-infoz-gz--fmad=falsez--opt-levelÚ0ú z--regAllocOptLevel=2z-vz--gpu-name=z-oT)ÚcheckÚ	close_fdsÚstderréÿ   z!Internal Triton PTX codegen erroré€   z`ptxas` raised SIGSEGVz`ptxas` failed with error code z
`ptxas` stderr:
z
Repro command: ú
zC

================================================================
z

zy
================================================================
please share the reproducer above with Triton project.
rV   )!r+   r   r&   r2   ÚtempfileÚNamedTemporaryFileÚwriteÚflushr    r
   rS  rT  r	   Údisable_ptxas_optrr   re   rp   rD   r0   r  Údump_ptxas_logrW   rr  rZ   ÚosÚexistsÚremoveÚCalledProcessErrorÚ
returncodeÚsignalÚSIGSEGVr¥   r   )r”   r9  râ   r  r^   r*   ÚfsrcÚflogÚfbinÚ
debug_infoÚfmadr&   Údisable_optÚptx_extra_optionsÚ	reg_allocÚ	ptxas_cmdÚlog_fileÚeÚlogÚerrorr\   rÅ   s                         r!   Ú
make_cubinzCUDABackend.make_cubinÒ  sp  € Ý˜$œ+Ô*Ñ+Ô+Ô0ˆÝÔ(°¸CÈÐOÑOÔOð J	 ÐSWÝÔ'¨u¸3ÀvÐNÑNÔNðJ	 ØRVØ�JŠJ�s‰OŒOˆOØ�JŠJ‰LŒLˆLØ”9˜tÑ#ˆDàˆJÝÔ Ô2ð ,à˜{Ð,BÐCÑC�
�
Ý”Ô/ð ,à˜t˜fÑ$�
�
ð ˜{˜mÑ+�
àÔ-ÐC�2�2°NÐ3CˆDÝ*¨:Ñ6Ô6ˆDõ 38´,Ô2PÐX˜=¨#Ð.Ð.ÐVXˆKð ?B¼oÐ U ¤× 5Ò 5°cÑ :Ô :Ð :ÐSUÐð 0Ð0ˆIð ðØ"ðØ%)ðØ+/ðØ2=ðØ@IðØL]ðØ_sÐmqÐ_sÐ_sðà”	ðàðà!%ðˆIð$(Ý”˜y°ÀÈdÐSÑSÔSÐSÝ”<Ô.ð /Ý˜dœi™œð /¨HÝ˜hŸmšm™oœoÑ.Ô.Ð.ð/ð /ð /ñ /ô /ð /ð /ð /ð /ð /ð /øøøð /ð /ð /ð /õ ”7—>’> $¤)Ñ,Ô,ð )Ý”I˜dœiÑ(Ô(Ð(Ý”7—>’> $¤)Ñ,Ô,ð )Ý”I˜dœiÑ(Ô(Ð(ùøÝÔ0ð (ð (ð (Ý˜$œ)‘_”_ð *¨Ø"Ÿ-š-™/œ/�Cð*ð *ð *ñ *ô *ð *ð *ð *ð *ð *ð *øøøð *ð *ð *ð *å”7—>’> $¤)Ñ,Ô,ð )Ý”I˜dœiÑ(Ô(Ð(à”< 3Ò&Ð&Ø?�E�EØ”\ S­6¬>Ñ%9Ò9Ð9Ø4�E�EàL¸a¼lÐLÐL�Eà!ð Cð CØ-0ðCð Cà+.¯8ª8°IÑ+>Ô+>ðCð Cð C�õ ð ð ðð ð
 ðð ð ñ ô ð õ ! Ñ'Ô'Ð'øøøøð5(øøøõ8 �d˜DÑ!Ô!ð ! QØŸš™œ�ð!ð !ð !ñ !ô !ð !ð !ð !ð !ð !ð !øøøð !ð !ð !ð !åŒw�~Š~˜dÑ#Ô#ð  Ý”	˜$‘”�ðUJ	 ð J	 ð J	 ñ J	 ô J	 ð J	 ð J	 ð J	 ð J	 ð J	 ð J	 øøøð J	 ð J	 ð J	 ð J	 ð J	 ð J	 ð J	 ñ J	 ô J	 ð J	 ð J	 ð J	 ð J	 ð J	 ð J	 øøøð J	 ð J	 ð J	 ð J	 ðV ˆsÈ   ¶OÁC$N*Ä3=H%Å0"FÆH%ÆF"Æ"H%Æ%F"Æ&A=H%È#N*È%L,È4L'ÉI)	ÉL'É)I-É-L'É0I-É1B6L'Ì'L,Ì,N*Ì?M ÍN*Í M$Í$N*Í'M$Í(6N*ÎOÎ*N.	Î.OÎ1N.	Î2OÏOÏOc                 ób  ‡ ‡‡— ‰                       ‰j        ¦  «        Š|t          j        k    rˆˆˆ fd„|d<   ˆˆˆ fd„|d<   n|t          j        k    r
ˆˆˆ fd„|d<   ˆˆˆ fd„|d<   ˆˆ fd„|d	<   ˆˆ fd
„|d<   t
          j        j        �%t
          j                             ‰ |‰|‰¦  «         d S d S )Nc                 ó4   •— ‰                      | |‰‰¦  «        S rL   )rô   ©r9  râ   r^   rN   r”   s     €€€r!   ú<lambda>z(CUDABackend.add_stages.<locals>.<lambda>$  s   ø€ °4·>²>À#ÀxÐQXÐZdÑ3eÔ3e€ r#   rú   c                 ó4   •— ‰                      | |‰‰¦  «        S rL   )r  rž  s     €€€r!   rŸ  z(CUDABackend.add_stages.<locals>.<lambda>%  s   ø€ °D·O²OÀCÈÐSZÐ\fÑ4gÔ4g€ r#   Úttgirc                 ó4   •— ‰                      | |‰‰¦  «        S rL   )r5  rž  s     €€€r!   rŸ  z(CUDABackend.add_stages.<locals>.<lambda>'  s   ø€ °D×4GÒ4GÈÈXÐW^Ð`jÑ4kÔ4k€ r#   c                 ó4   •— ‰                      | |‰‰¦  «        S rL   )r=  rž  s     €€€r!   rŸ  z(CUDABackend.add_stages.<locals>.<lambda>(  s   ø€ ¨t¯~ª~¸cÀ8ÈWÐV`Ñ/aÔ/a€ r#   Úllirc                 óH   •— ‰                      | |‰‰j        j        ¦  «        S rL   )rt  r   r&   ©r9  râ   rN   r”   s     €€r!   rŸ  z(CUDABackend.add_stages.<locals>.<lambda>)  s    ø€ ¨d¯mªm¸CÀÈ7ÐTXÔT_ÔTdÑ.eÔ.e€ r#   Úptxc                 óH   •— ‰                      | |‰‰j        j        ¦  «        S rL   )r›  r   r&   r¦  s     €€r!   rŸ  z(CUDABackend.add_stages.<locals>.<lambda>*  s    ø€ °·²ÀÀXÈwÐX\ÔXcÔXhÑ0iÔ0i€ r#   rÅ   )rÀ   r&   r   ÚTRITONÚGLUONr
   rÕ   Úadd_stages_inspection_hook)r”   ÚstagesrN   rÜ   r^   s   ` ` @r!   Ú
add_stageszCUDABackend.add_stages!  sù   øøø€ Ø×%Ò% g¤lÑ3Ô3ˆ
Ø•x”Ò&Ð&ØeÐeÐeÐeÐeÐeˆF�6‰NØgÐgÐgÐgÐgÐgˆF�7‰OˆOØ�œÒ'Ð'ØkÐkÐkÐkÐkÐkˆF�7‰OØaÐaÐaÐaÐaÐaˆˆv‰ØeÐeÐeÐeÐeˆˆu‰ØiÐiÐiÐiÐiˆˆw‰ÝŒ=Ô3Ð?ÝŒM×4Ò4°T¸6À7ÈHÐV`ÑaÔaÐaÐaÐað @Ð?r#   c                 óV   — t          | j        j        ¦  «        }|› d| j        j        › �S )NrŸ   )r6   r   r&   )r”   r5   s     r!   r©   zCUDABackend.hash.  s-   € å# D¤KÔ$4Ñ5Ô5ˆØÐ.Ð.˜DœKÔ,Ð.Ð.Ð.r#   )r­   r®   r¯   rð   Ústaticmethodr   r¸   rÀ   rB   rÃ   rÇ   r   rß   rã   rë   r   r   rí   rï   rô   r  r5  r=  rt  r›  r­  Ú	functoolsÚ	lru_cacher©   Ú__classcell__)rÉ   s   @r!   rµ   rµ   š   sÃ  ø€ € € € € Ø€Oàð( 	ð (ð (ð (ñ „\ð(ð#ð #ð #ð$¨#ð $ð $ð $ð $ð"˜yð "¨Tð "ð "ð "ð "ð "ð "ð# Sð #ð #ð #ð #ð@
ð 
ð 
ðð ð ð>  S¨* _Ô 5ð >ð >ð >ð >ð;ð ;ð ;ð
 ðð ñ „\ðð  ðGð Gñ „\ðGðRð ð ð&^ð ^ð ^ð@ð ð ð4Mð Mð Mð^bð bð bð €YÔÑÔð/ð /ñ Ôð/ð /ð /ð /ð /r#   rµ   )r,   ),Útriton.backends.compilerr   r   r   Útriton._C.libtritonr   r   r   r	   Útritonr
   Útriton.runtime.errorsr   Údataclassesr   r°  Útypingr   r   r   r   Útypesr   rX   rº   r�  rŒ  r‡  r0   Úpathlibr   r%   r$   Ú
NvidiaToolr+   r±  r6   rJ   rO   rT   r]   re   rh   rµ   rc   r#   r!   ú<module>r¼     sÓ  ðØ EÐ EÐ EÐ EÐ EÐ EÐ EÐ EÐ EÐ EØ 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ø Ð Ð Ð Ð Ð Ø ,Ð ,Ð ,Ð ,Ð ,Ð ,à !Ð !Ð !Ð !Ð !Ð !Ø Ð Ð Ð Ø -Ð -Ð -Ð -Ð -Ð -Ð -Ð -Ð -Ð -Ð -Ð -Ø Ð Ð Ð Ð Ð Ø €€€Ø 	€	€	€	Ø €€€Ø €€€Ø 	€	€	€	Ø Ð Ð Ð Ø Ð Ð Ð Ð Ð ð#˜ð #ð #ð #ð #ðO�Cð O˜EÔ,ð Oð Oð Oð Oð €ÔÑÔðð ˜Cð ð ð ñ Ôðð €ÔÑÔði Sð ið ið iñ Ôðið.°ð ð ð ð ð €ÔÑÔð
 ð 
ð 
ð 
ñ Ôð
ð €Ô�TÑÔð4ð 4ñ Ôð4ð
&¨ð &ð &ð &ð &ð €�$ÐÑÔð-4ð -4ð -4ð -4ð -4ñ -4ô -4ñ Ôð-4ð`W/ð W/ð W/ð W/ð W/�+ñ W/ô W/ð W/ð W/ð W/r#   