§
    8˜ºh
K  ã                   ó   — 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
j         fd„Z! ej"        ¦   «         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   r   )ÚscalarÚprimitive_bitwidth)Úlhs_typeÚrhs_typeÚlhs_bitwidthÚrhs_bitwidths       úg/var/www/html/services/rag-service/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"   r   c                  ó$   — t           j        j        S ©N)r
   r	   Úptxas© r"   r    Ú	get_ptxasr)   !   s   € ÝŒ<ÔÐr"   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)Úmock_verÚversions     r    Úget_ptxas_versionr3   %   sF   € åŒ|Ô,€HØÐØˆÝÔ%¥y¡{¤{Ô'7¸Ð&EÑFÔF×MÒMÈgÑVÔV€GØ€Nr"   c                 ó  — 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 t          d
| z   ¦  «        ‚)zK
    Get the highest PTX version supported by the current CUDA driver.
    ú.é   é   éP   é   é   éF   é
   é?   z?Triton only support CUDA 10.0 or higher, but got CUDA version: )Ú
isinstanceÚstrÚmapr#   ÚsplitÚRuntimeError)Úcuda_versionÚmajorÚminors      r    Úptx_get_versionrF   .   s    € õ
 �l¥CÑ(Ô(Ð(Ð(Ð(Ý•s˜L×.Ò.¨sÑ3Ô3Ñ4Ô4�L€Eˆ5Ø�‚{€{Ø�1Š9ˆ9Ø˜‘:Ðà˜‘: ‘>Ð!Ø�‚{€{Ø�E‰zÐØ�‚{€{Ø�E‰zÐÝ
ÐXÐ[gÑgÑ
hÔ
hÐhr"   Úarchc                 ó\   — | j         }|€"t          ¦   «         j        }t          |¦  «        }|S r&   )Úptx_versionr)   r2   rF   )ÚoptionsrG   rI   rC   s       r    Úget_ptx_version_from_optionsrK   A   s/   € ØÔ%€KØÐÝ ‘{”{Ô*ˆÝ% lÑ3Ô3ˆØÐr"   c                 óP   — t          | |¦  «        }t          d|¦  «        }d|› �}|S )NéV   z+ptx)rK   Úmin)rJ   rG   rI   Úllvm_ptx_versionÚfeaturess        r    Úget_featuresrQ   I   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)r/   Úfs     r    Ú	file_hashrZ   W   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 )NéZ   ÚaÚ Úsm_r(   )r[   Úsuffixs     r    Úsm_arch_from_capabilityrb   ]   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         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<   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 d#„ Z!dS )$ÚCUDAOptionsé   Ú	num_warpsr9   Únum_ctasé   Ú
num_stagesNÚmaxnreg)r9   r9   r9   Úcluster_dimsrI   Úptx_optionsÚir_overrideTÚenable_fp_fusionFÚlaunch_cooperative_gridÚ
launch_pdl)Úfp8e5Úfp8e4b15Úsupported_fp8_dtypesr(   Ú!deprecated_fp8_dot_operand_dtypesÚtf32Údefault_dot_input_precision)rv   Útf32x3ÚieeeÚallowed_dot_input_precisionsÚmax_num_imprecise_acc_defaultÚextern_libsÚdebugÚcudaÚbackend_nameÚsanitize_overflowrG   c                 ó¼  — 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   r9   znum_warps must be a power of 2)r   Ú__file__Úparentr|   ÚdictÚgetr
   r	   Úlibdevice_pathr?   ÚobjectÚ__setattr__ÚtupleÚitemsrg   )Ú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 r&   )rZ   )Ú.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 )ú-r(   )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Œ   rU   rV   ÚencoderX   )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"   )"Ú__name__Ú
__module__Ú__qualname__rg   r#   Ú__annotations__rh   rj   rk   r   rl   r‹   rI   rm   r?   rn   ro   Úboolrp   rq   rt   r   ru   rw   rz   r{   r|   r†   r}   r   r€   rG   r�   r¢   r(   r"   r    re   re   c   s¢  € € € € € € à€IˆsÐÐÑØ€HˆcÐÐÑØ€J�ÐÐÑð "€GˆX�cŒ]Ð!Ð!Ñ!Ø#€L�%Ð#Ð#Ñ#Ø€K�ÐÐÑØ€K�ÐÐÑØ!%€K�˜#”Ð%Ð%Ñ%Ø!Ð�dÐ!Ð!Ñ!Ø$)Ð˜TÐ)Ð)Ñ)Ø€J�ÐÐÑØ'<Ð˜% œ*Ð<Ð<Ñ<Ø46Ð% u¨S¤zÐ6Ð6Ñ6Ø'-Ð Ð-Ð-Ñ-Ø/IÐ  %¨¤*ÐIÐIÑIØ*.Ð! 4Ð.Ð.Ñ.Ø€K�ÐÐÑØ€Eˆ4ÐÐÑØ€L�#ÐÐÑØ"Ð�tÐ"Ð"Ñ"Ø€Dˆ#ÐÐÑð0ð 0ð 0ð?ð ?ð ?ð ?ð ?r"   re   c                   ó
  ‡ — e 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 )ÚCUDABackendr   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 r9   )ÚreÚ	fullmatchÚ
ValueErrorr#   Úgroup)r�   rG   Ú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´   rG   )r�   rJ   r[   s      r    Úget_target_namezCUDABackend.get_target_name›   s%   € Ø×%Ò% g¤lÑ3Ô3ˆ
Ø#�zÐ#Ð#Ð#r"   Nc                 óX   •— t          ¦   «                              |¦  «         d| _        d S )NÚcubin)ÚsuperÚ__init__Ú
binary_ext)r�   r   Ú	__class__s     €r    r»   zCUDABackend.__init__Ÿ   s&   ø€ Ý‰Œ×Ò˜Ñ Ô Ð Ø!ˆŒˆˆr"   c                 ó@  ‡— dt           j        j        pd| j        j        › �i}|                     ˆfd„t          j                             ¦   «         D ¦   «         ¦  «         t          |  
                    |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 )NrG   Úsmc                 ó:   •— i | ]}|‰v ¯‰|         ®|‰|         “ŒS r&   r(   )r’   r“   Úoptss     €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"   rt   éY   Úfp8e4nvru   r]   )rs   ro   i   @r   r{   r(   )r
   ÚruntimeÚoverride_archr   rG   Úupdatere   Ú__dataclass_fields__Úkeysr#   r´   Úsetrt   Úaddr‹   r�   ÚlanguageÚdefault_fp_fusion)r�   rÁ   Úargsr[   rt   s    `   r    Úparse_optionszCUDABackend.parse_options£   s:  ø€ Ø�œÔ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ˆ
à!¨Ð-Ð-Ý#&¥{Ô'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                 ór   — |j         |j        |j        |j        d         |j        d         |j        d         fS )Nr   r9   é   )rg   rh   Úsharedrl   )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   r8   )Úconvert_custom_typesr$   )Útriton.language.extra.cudarÌ   Úextrar~   r#   r´   rG   Úconvert_custom_float8_sm80Úconvert_custom_float8_sm70r$   r   )r�   rJ   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        |¦  «         d S r&   )r	   Úload_dialects)r�   Úctxs     r    rà   zCUDABackend.load_dialectsÑ   s   € ÝÔ˜SÑ!Ô!Ð!Ð!Ð!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                             |¦  «         |                     | ¦  «         | S )Nr<   é	   )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    Ú	make_ttirzCUDABackend.make_ttirÔ   s  € åŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐÝŒ×!Ò! "Ñ%Ô%Ð%ÝŒ×.Ò.¨rÑ2Ô2Ð2Ø˜Ñ˜aÒÐÝŒK×@Ò@ÀÑDÔDÐDÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò Ñ#Ô#Ð#ÝŒ×)Ò)¨"Ñ-Ô-Ð-ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒ×$Ò$ RÑ(Ô(Ð(ÝŒ×#Ò# BÑ'Ô'Ð'Ø
�Šˆs‰ŒˆØˆ
r"   c                 ó.  — |j         �E|                      dt          j        | j        ¦  «                             |j         ¦  «        ¦  «         t          j        ¦   «         }|j        �6|j        d         |_	        |j        d         |_
        |j        d         |_        t          j        | j        ¦  «        }|                     ¦   «         }t          j                             |d|› �|j        d|j        ¦  «         t          j                             |¦  «         |dz  dk    r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         *                    |¦  «         t          j        j         +                    |¦  «         t          j         &                    ||j%        ¦  «         t          j         '                    |¦  «         t          j         ,                    ||j%        ¦  «         t          j         (                    ||j%        |¦  «         t          j         "                    |¦  «         t          j        j         -                    |¦  «         nt          j         !                    |¦  «         t          j                              |¦  «         t          j                             |¦  «         t          j         .                    |¦  «         t          j                             ||d	k    ¦  «         t          j         /                    |¦  «         t          j        j         0                    |¦  «         t          j                             |¦  «         t          j        j         1                    |¦  «         t          j         2                    |¦  «         t          j         3                    |¦  «         t          j                             |¦  «         t          j         4                    |¦  «         |dz  dk    rHt          j        j         5                    |¦  «         t          j        j         6                    |¦  «         t          j         7                    |¦  «         t          j                              |¦  «         | 8                    | ¦  «         |j	        |j
        |j        f|d<   |  9                    ¦   «         }||d<   | S )Nzttg.maxnregr   r9   rÑ   r¶   r   r<   r   r8   )r   rã   rã   rl   Útensordesc_meta):rk   Úset_attrr   Úbuilderrå   Úget_int32_attrr	   ÚClusterInforl   ÚclusterDimXÚclusterDimYÚclusterDimZrä   ræ   r   ré   Úadd_convert_to_ttgpuirrg   rh   ÚttgpuirÚadd_coalesceÚadd_f32_dot_tcÚ	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_warpspecrj   Ú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_remove_tmem_tokensÚadd_prefetchÚadd_coalesce_async_copyÚadd_optimize_tmem_layoutsÚadd_interleave_tmemÚadd_reduce_data_duplicationÚadd_reorder_instructionsrð   Úadd_tma_loweringÚadd_fence_insertionÚadd_sccprò   Úget_tensordesc_metadata)ró   rÓ   rô   r[   Úcluster_inforõ   Údump_enabledrø   s           r    Ú
make_ttgirzCUDABackend.make_ttgirå   sË  € ð Œ;Ð"Ø�LŠL˜­¬
°3´;Ñ(?Ô(?×(NÒ(NÈsÌ{Ñ([Ô([Ñ\Ô\Ð\åÔ)Ñ+Ô+ˆØÔÐ'Ø'*Ô'7¸Ô':ˆLÔ$Ø'*Ô'7¸Ô':ˆLÔ$Ø'*Ô'7¸Ô':ˆLÔ$ÝŒ_˜Sœ[Ñ)Ô)ˆØ—’Ñ(Ô(ˆÝŒ×*Ò*¨2Ð/C°zÐ/CÐ/CÀSÄ]ÐTVÐX[ÔXdÑeÔeÐeåŒ×#Ò# BÑ'Ô'Ð'Ø˜Ñ˜qÒ Ð ÝŒN×)Ò)¨"Ñ-Ô-Ð-åŒÔ×,Ò,¨R°Ñ>Ô>Ð>ÝŒ×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×/Ò/°Ñ3Ô3Ð3ÝŒ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×;Ò;¸BÑ?Ô?Ð?ÝŒMÔ#×:Ò:¸2Ñ>Ô>Ð>Ð>åŒK×'Ò'¨Ñ+Ô+Ð+ÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×&Ò& rÑ*Ô*Ð*ÝŒ×#Ò# BÑ'Ô'Ð'ÝŒ×0Ò0°°ZÀ2Ò5EÑFÔFÐFÝŒ×.Ò.¨rÑ2Ô2Ð2ÝŒÔ×9Ò9¸"Ñ=Ô=Ð=ÝŒ×4Ò4°RÑ8Ô8Ð8ÝŒÔ×3Ò3°BÑ7Ô7Ð7ÝŒ×2Ò2°2Ñ6Ô6Ð6ÝŒ×/Ò/°Ñ3Ô3Ð3ÝŒ×&Ò& rÑ*Ô*Ð*ÝŒ×$Ò$ RÑ(Ô(Ð(Ø˜Ñ˜qÒ Ð ÝŒMÔ#×4Ò4°RÑ8Ô8Ð8ÝŒMÔ#×7Ò7¸Ñ;Ô;Ð;ÝŒ×Ò˜rÑ"Ô"Ð"ÝŒ×'Ò'¨Ñ+Ô+Ð+Ø
�Šˆs‰ŒˆØ$0Ô$<¸lÔ>VÐXdÔXpÐ#qˆ�Ñ Ø×5Ò5Ñ7Ô7ˆØ&5ˆÐ"Ñ#Øˆ
r"   c                 óò  — |}t          j        |j        ¦  «        }|                     ¦   «          t          j                             |¦  «         t          j                             |¦  «         t          j	         
                    |¦  «         t          j                             |¦  «         t          j                             |¦  «         |                     |¦  «         |                     ¦   «         |d<   |S )Nrø   )r   rä   rå   ræ   r   r  rè   rç   r!  ré   r  rì   r  rò   r"  )r�   ÚsrcrÓ   rJ   r[   ró   rõ   s          r    Ú	ttgir_optzCUDABackend.ttgir_opt0  sÁ   € ØˆÝŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐåŒ×"Ò" 2Ñ&Ô&Ð&ÝŒ×Ò˜rÑ"Ô"Ð"ÝŒ×&Ò& rÑ*Ô*Ð*ÝŒ×(Ò(¨Ñ,Ô,Ð,ÝŒ×7Ò7¸Ñ;Ô;Ð;à
�Šˆs‰ŒˆØ&)×&AÒ&AÑ&CÔ&CˆÐ"Ñ#Øˆ
r"   c                 ó4  — t          || j        j        ¦  «        }|}t          j        |j        ¦  «        }|                     ¦   «          t          j        j	         
                    |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j        j	                             |¦  «         t          j                             |¦  «         t          j        j                             |||¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j        j	                             |¦  «         t          j        j	                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t4          j        j        st          j                             |¦  «         |                     |¦  «         tA          j!        ¦   «          tA          j        ¦   «         }t4          j        j"        rtG          d¦  «        ‚tA          j$        ||¦  «        }	tK          |¦  «        }
tM          || j        j        ¦  «        }d}t          j'        ¦   «          tA          j(        |	||
|¦  «         t          j)        |	¦  «         |j*        r&d„ |j*        D ¦   «         }tA          j+        |	|¦  «         tA          j,        |	t@          j-        ¦  «         | .                    d¦  «        }|�||d<   | .                    d¦  «        |d<   | .                    d¦  «        |d	<   | .                    d
¦  «        |d<   | .                    d¦  «        |d<   t_          |	¦  «        }~	~|S )NzYAddress Sanitizer Error: Address sanitizer is currently only supported on the AMD backendúnvptx64-nvidia-cudac                 ó   — g | ]\  }}|‘ŒS r(   r(   )r’   r™   r/   s      r    r›   z)CUDABackend.make_llir.<locals>.<listcomp>h  s   € ÐBÐBÐB™l˜t T�TÐBÐBÐBr"   zttg.total-num-warpsrg   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_align)0rK   r   rG   r   rä   rå   ræ   r	   r   r  Úadd_lower_mmar  r  Úadd_allocate_warp_groupsÚconvertÚadd_scf_to_cfÚadd_allocate_shared_memoryÚadd_allocate_tensor_memoryÚ"add_allocate_global_scratch_memoryÚadd_to_llvmirrç   rì   rï   Úadd_nvgpu_to_llvmÚadd_warp_specialize_to_llvmrð   r
   ÚcompilationÚdisable_line_infoÚllvmirÚadd_di_scoperò   r   Úinit_targetsÚenable_asanrB   Ú	to_modulerb   rQ   Úset_short_ptrÚattach_datalayoutÚset_nvvm_reflect_ftzr|   Úlink_extern_libsÚoptimize_moduleÚOPTIMIZE_O3Úget_int_attrr?   )r�   r'  rÓ   rJ   r[   rI   ró   rõ   rå   Úllvm_modÚprocrP   ÚtripleÚpathsÚtotal_num_warpsÚrets                   r    Ú	make_llirzCUDABackend.make_llir?  sf  € Ý2°7¸D¼KÔ<LÑMÔMˆàˆåŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐåŒÔ×-Ò-¨bÑ1Ô1Ð1ÝŒ×7Ò7¸Ñ;Ô;Ð;ÝŒ×/Ò/°Ñ3Ô3Ð3ÝŒ×$Ò$ RÑ(Ô(Ð(ÝŒ×1Ò1°"Ñ5Ô5Ð5ÝŒÔ×:Ò:¸2Ñ>Ô>Ð>ÝŒ×9Ò9¸"Ñ=Ô=Ð=ÝŒÔ×+Ò+¨B°
¸KÑHÔHÐHÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒÔ×1Ò1°"Ñ5Ô5Ð5ÝŒÔ×;Ò;¸BÑ?Ô?Ð?ÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒ×$Ò$ RÑ(Ô(Ð(ÝÔ Ô2ð 	+ÝŒM×&Ò& rÑ*Ô*Ð*Ø
�Šˆs‰ŒˆåÔÑÔÐÝ”,‘.”.ˆÝÔÔ(ð 	mÝØkñmô mð må”> # wÑ/Ô/ˆÝ& zÑ2Ô2ˆÝ ¨¬Ô)9Ñ:Ô:ˆØ&ˆÝÔÑÔÐÝÔ˜x¨°°xÑ@Ô@Ð@ÝÔ# HÑ-Ô-Ð-àÔð 	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ˆÐ'Ñ(Ý�(‰mŒmˆØØØˆ
r"   c           	      óp  — t          || j        j        ¦  «        }d}t          |¦  «        }t	          || j        j        ¦  «        }t          j        ||||g |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        dd|	¦  «        }	t          j        j        rt!          d¦  «         t!          |	¦  «         |	S )Nr*  Fz(.visible .entry ([a-zA-Z_][a-zA-Z0-9_]*)r9   r   r™   r<   r5   z\.version \d+\.\d+z	.version )Úflagsz\.target sm_\d+z.target sm_z,\s*debug|debug,\s*r_   z // -----// NVPTX Dump //----- //)rK   r   rG   rb   rQ   r   Útranslate_to_asmro   r®   ÚfindallÚlenÚsubÚ	MULTILINEr
   r	   Ú
dump_nvptxÚprint)r�   r'  rÓ   rô   r[   rI   rI  rH  rP   rL  Únamess              r    Úmake_ptxzCUDABackend.make_ptx{  s6  € Ý2°3¸¼Ô8HÑIÔIˆà&ˆÝ& zÑ2Ô2ˆÝ  T¤[Ô%5Ñ6Ô6ˆÝÔ# C¨°°xÀÀSÔEYÐ[`ÑaÔaˆå”
Ð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\Ð]Ñ]Ô]ˆåŒfÐ+¨R°Ñ5Ô5ˆÝŒ<Ô"ð 	ÝÐ4Ñ5Ô5Ð5Ý�#‰JŒJˆJØˆ
r"   c                 óL  — t          ¦   «         j        }t          j        ddd¬¦  «        5 }t          j        ddd¬¦  «        5 }|                     |¦  «         |                     ¦   «          |j        dz   }t          j        j	        rdd	gndg}	|j
        rg nd
g}
t          |¦  «        }t          j        j        rddgng }|j        r|j                             d¦  «        ng }|g|	¢|
¢d‘|¢|¢d|› �‘|j        ‘d‘|‘}	 t!          j        |dd|¬¦  «         t$          j                             |j        ¦  «        rt%          j        |j        ¦  «         t$          j                             |j        ¦  «        rt%          j        |j        ¦  «         nò# t           j        $ rà}t-          |j        ¦  «        5 }|                     ¦   «         }d d d ¦  «         n# 1 swxY w Y   t$          j                             |j        ¦  «        rt%          j        |j        ¦  «         |j        dk    rd}n%|j        dt2          j        z   k    rd}n
d|j        › �}t7          |› d|› dd                     |¦  «        › d�¦  «        ‚d }~ww xY wt-          |d¦  «        5 }|                     ¦   «         }d d d ¦  «         n# 1 swxY w Y   t$          j                             |¦  «        rt%          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)ÚdeleteÚmodera   Úrz.logz.oz	-lineinfoz-suppress-debug-infoz--fmad=falsez--opt-levelÚ0ú z-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: ú
rS   )r)   r/   ÚtempfileÚNamedTemporaryFileÚwriteÚflushr™   r
   r9  r:  ro   rb   r	   Údisable_ptxas_optrm   rA   r-   rò   ÚosÚexistsÚremoveÚCalledProcessErrorrT   rW   Ú
returncodeÚsignalÚSIGSEGVr   rž   )r�   r'  rÓ   rô   r[   r'   ÚfsrcÚflogÚfbinÚ	line_infoÚfmadrG   Údisable_optÚptx_extra_optionsÚ	ptxas_cmdÚeÚlog_fileÚlogÚerrorrY   r¹   s                        r    Ú
make_cubinzCUDABackend.make_cubin‘  sM  € Ý‘”Ô ˆÝÔ(°¸CÈÐOÑOÔOð .	 ÐSWÝÔ'¨u¸3ÀvÐNÑNÔNð.	 ØRVØ�JŠJ�s‰OŒOˆOØ�JŠJ‰LŒLˆLØ”9˜tÑ#ˆDåAFÔARÔAdÐw˜Ð&<Ð=Ð=ÐkvÐjwˆIØÔ-ÐC�2�2°NÐ3CˆDÝ*¨:Ñ6Ô6ˆDõ 38´,Ô2PÐX˜=¨#Ð.Ð.ÐVXˆKð ?B¼oÐ U ¤× 5Ò 5°cÑ :Ô :Ð :ÐSUÐð ðØ!ðØ$(ðØ*.ðØ1<ðØ?PðØRfÐ`dÐRfÐRfðØhlÔhqðØswðàðˆIðLÝ”˜y°ÀÈdÐSÑSÔSÐSÝ”7—>’> $¤)Ñ,Ô,ð )Ý”I˜dœiÑ(Ô(Ð(Ý”7—>’> $¤)Ñ,Ô,ð )Ý”I˜dœiÑ(Ô(Ð(øøÝÔ0ð Lð Lð LÝ˜$œ)‘_”_ð *¨Ø"Ÿ-š-™/œ/�Cð*ð *ð *ñ *ô *ð *ð *ð *ð *ð *ð *øøøð *ð *ð *ð *å”7—>’> $¤)Ñ,Ô,ð )Ý”I˜dœiÑ(Ô(Ð(à”< 3Ò&Ð&Ø?�E�EØ”\ S­6¬>Ñ%9Ò9Ð9Ø4�E�EàL¸a¼lÐLÐL�Eå  Eð "Kð "KØ58ð"Kð "Kà36·8²8¸IÑ3FÔ3Fð"Kð "Kð "Kñ Lô Lð LøøøøðLøøøõ" �d˜DÑ!Ô!ð ! QØŸš™œ�ð!ð !ð !ñ !ô !ð !ð !ð !ð !ð !ð !øøøð !ð !ð !ð !åŒw�~Š~˜dÑ#Ô#ð  Ý”	˜$‘”�ð].	 ð .	 ð .	 ñ .	 ô .	 ð .	 ð .	 ð .	 ð .	 ð .	 ð .	 øøøð .	 ð .	 ð .	 ð .	 ð .	 ð .	 ð .	 ñ .	 ô .	 ð .	 ð .	 ð .	 ð .	 ð .	 ð .	 øøøð .	 ð .	 ð .	 ð .	 ð^ ˆs¤   «LÁB>LÄBFÆLÆJÆ$I?Æ8G	ÇI?ÇGÇI?Ç GÇ!BI?É?JÊLÊJ8Ê,LÊ8J<Ê<LÊ?J<Ë 6LË6LÌL	ÌLÌ	L	Ì
LÌLÌ Lc                 óö   ‡ ‡‡— ‰                       ‰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<   d S )Nc                 ó4   •— ‰                      | |‰‰¦  «        S r&   )rö   ©r'  rÓ   r[   rJ   r�   s     €€€r    ú<lambda>z(CUDABackend.add_stages.<locals>.<lambda>Ç  s   ø€ °4·>²>À#ÀxÐQXÐZdÑ3eÔ3e€ r"   ré   c                 ó4   •— ‰                      | |‰‰¦  «        S r&   )r%  r�  s     €€€r    r‚  z(CUDABackend.add_stages.<locals>.<lambda>È  s   ø€ °D·O²OÀCÈÐSZÐ\fÑ4gÔ4g€ r"   Úttgirc                 ó4   •— ‰                      | |‰‰¦  «        S r&   )r(  r�  s     €€€r    r‚  z(CUDABackend.add_stages.<locals>.<lambda>Ê  s   ø€ °D·N²NÀ3ÈÐRYÐ[eÑ4fÔ4f€ r"   c                 ó4   •— ‰                      | |‰‰¦  «        S r&   )rM  r�  s     €€€r    r‚  z(CUDABackend.add_stages.<locals>.<lambda>Ë  s   ø€ ¨t¯~ª~¸cÀ8ÈWÐV`Ñ/aÔ/a€ r"   Úllirc                 óH   •— ‰                      | |‰‰j        j        ¦  «        S r&   )rX  r   rG   ©r'  rÓ   rJ   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 r&   )r~  r   rG   r‰  s     €€r    r‚  z(CUDABackend.add_stages.<locals>.<lambda>Í  s    ø€ °·²ÀÀXÈwÐX\ÔXcÔXhÑ0iÔ0i€ r"   r¹   )r´   rG   r   ÚTRITONÚGLUON)r�   ÚstagesrJ   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Ø�œÒ'Ð'ØfÐfÐfÐfÐfÐfˆF�7‰OØaÐaÐaÐaÐaÐaˆˆv‰ØeÐeÐeÐeÐeˆˆu‰ØiÐiÐiÐiÐiˆˆw‰ˆˆr"   c                 ó@   — t          ¦   «         }|› d| j        j        › �S )Nr˜   )r3   r   rG   )r�   r2   s     r    r¢   zCUDABackend.hashÏ  s&   € å#Ñ%Ô%ˆØÐ.Ð.˜DœKÔ,Ð.Ð.Ð.r"   )r£   r¤   r¥   Ústaticmethodr   r¬   r´   r?   r·   r»   r   rÏ   rÔ   rÜ   r   r   rÞ   rà   rö   r%  r(  rM  rX  r~  r�  Ú	functoolsÚ	lru_cacher¢   Ú__classcell__)r½   s   @r    r©   r©   Ž   s·  ø€ € € € € àð( 	ð (ð (ð (ñ „\ð(ð#ð #ð #ð$¨#ð $ð $ð $ð $ð"˜yð "¨Tð "ð "ð "ð "ð "ð "ð# Sð #ð #ð #ð #ð,
ð 
ð 
ðð ð ð>  S¨* _Ô 5ð >ð >ð >ð >ð"ð "ð "ð ðð ñ „\ðð  ðHð Hñ „\ðHðTð ð ð:ð :ð :ðxð ð ð,1ð 1ð 1ðf	jð 	jð 	jð €YÔÑÔð/ð /ñ Ôð/ð /ð /ð /ð /r"   r©   ),Útriton.backends.compilerr   r   r   Útriton._C.libtritonr   r   r   r	   Útritonr
   Útriton.runtime.errorsr   Údataclassesr   r’  Útypingr   r   r   r   Útypesr   rU   r®   rf  rp  rk  r-   Úpathlibr   r$   Ú
NvidiaToolr)   r“  r3   r#   rF   rK   rQ   rZ   rb   re   r©   r(   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Ø Ð Ð Ð Ð Ð Ø ,Ð ,Ð ,Ð ,Ð ,Ð ,à !Ð !Ð !Ð !Ð !Ð !Ø Ð Ð Ð Ø -Ð -Ð -Ð -Ð -Ð -Ð -Ð -Ð -Ð -Ð -Ð -Ø Ð Ð Ð Ð Ð Ø €€€Ø 	€	€	€	Ø €€€Ø €€€Ø 	€	€	€	Ø Ð Ð Ð Ø Ð Ð Ð Ð Ð ð#˜ð #ð #ð #ð #ð�5Ô#ð ð ð ð ð €ÔÑÔðð ñ Ôðð €ÔÑÔði Sð ið ið iñ Ôðið$°ð ð ð ð ð €ÔÑÔð
 ð 
ð 
ð 
ñ Ôð
ð €Ô�TÑÔð4ð 4ñ Ôð4ð
&¨ð &ð &ð &ð &ð €�$ÐÑÔð'?ð '?ð '?ð '?ð '?ñ '?ô '?ñ Ôð'?ðTD/ð D/ð D/ð D/ð D/�+ñ D/ô D/ð D/ð D/ð D/r"   