Ë
    ÎÍ:j
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jD                  «       d„ «       Z# ejD                  «       de$fd„«       Z%de$fd„Z& ejD                  «       de$fd„«       Z' ejD                  d«      d„ «       Z(de$fd„Z) ed¬«       G d„ d«      «       Z* G d„ de«      Z+y)é    )ÚBaseBackendÚ	GPUTargetÚLanguage)ÚirÚpassesÚllvmÚnvidia)Úknobs)Ú
PTXASError)Ú	dataclassN)ÚAnyÚDictÚTupleÚOptional)Ú
ModuleType)ÚPathÚtargetc                 ó>   — dt         t        t        t        f   fd„}|S )NÚreturnc                 ó€   — | j                   j                  }|j                   j                  }||k(  sJ d«       ‚|dk(  ryy)Nz%lhs and rhs bitwidth must be the sameé   )é   r   é    )r   r   r   )ÚscalarÚprimitive_bitwidth)Úlhs_typeÚrhs_typeÚlhs_bitwidthÚrhs_bitwidths       út/home/mcse/projects/srt_converter/srt-converter-venv/lib/python3.12/site-packages/triton/backends/nvidia/compiler.pyÚcheck_dot_compatibilityz-min_dot_size.<locals>.check_dot_compatibility   sB   € Ø—‘×9Ñ9ˆØ—‘×9Ñ9ˆØ˜|Ò+ÐTÐ-TÓTÐ+Ø˜1ÒØàó    )r   Úint)r   r!   s     r    Úmin_dot_sizer$      s!   € ð ´u¼SÄ#Äs¸]Ñ7Kó  ð #Ð"r"   r   c                  ó6   — 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«      j                  d«      }|S )Nz	--versionúutf-8)r
   r	   Úmock_ptx_versionÚ
subprocessÚcheck_outputr)   ÚpathÚdecode)Úmock_verÚversions     r    Úget_ptxas_versionr3   %   sI   € ä�|‰|×,Ñ,€HØÐØˆÜ×%Ñ%¤y£{×'7Ñ'7¸Ð&EÓF×MÑMÈgÓV€GØ€Nr"   c                 óÜ   — t        | t        «      sJ ‚t        t        | j	                  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Ó4�L€Eˆ5Ø�‚{Ø�1Š9Ø˜‘:Ðà˜‘: ‘>Ð!Ø�‚{Ø�E‰zÐØ�‚{Ø�E‰zÐÜ
ÐXÐ[gÑgÓ
hÐhr"   Úarchc                 ó`   — | j                   }|€t        «       j                  }t        |«      }|S r&   )Úptx_versionr)   r2   rF   )ÚoptionsrG   rI   rC   s       r    Úget_ptx_version_from_optionsrK   A   s0   € Ø×%Ñ%€KØÐÜ “{×*Ñ*ˆÜ% lÓ3ˆØÐr"   c                 ó@   — t        | |«      }t        d|«      }d|› �}|S )NéV   z+ptx)rK   Úmin)rJ   rG   rI   Úllvm_ptx_versionÚfeaturess        r    Úget_featuresrQ   I   s0   € ä.¨w¸Ó=€Kô ˜2˜{Ó+ÐØÐ&Ð'Ð(€HØ€Or"   c                 ó¨   — t        | d«      5 }t        j                  |j                  «       «      j	                  «       cd d d «       S # 1 sw Y   y xY w)NÚrb)ÚopenÚhashlibÚsha256ÚreadÚ	hexdigest)r/   Úfs     r    Ú	file_hashrZ   W   s>   € ä	ˆd�DÓ	ð 4˜QÜ�~‰~˜aŸf™f›hÓ'×1Ñ1Ó3÷4÷ 4ò 4ús   �1AÁAÚ
capabilityc                 ó"   — | dk\  rdnd}d| › |› �S )NéZ   ÚaÚ Úsm_r(   )r[   Úsuffixs     r    Úsm_arch_from_capabilityrb   ]   s!   € à "Ò$‰S¨"€FØ��˜V˜HÐ%Ð%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!y)$Ú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                  «      }|j                  dd «      s-t        j                  j                  xs t        |dz  «      |d<   t        j                  | dt        |j                  «       «      «       | j                  dkD  r| j                  | j                  dz
  z  dk(  sJ d«       ‚y )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¼$¸t×?OÑ?OÓ:PˆØ�‰˜{¨DÔ1Ü',§|¡|×'BÑ'BÒ'mÄcÈ.Ð[lÑJlÓFmˆK˜Ñ$ä×Ñ˜4 ´°k×6GÑ6GÓ6IÓ0JÔKØ�~‰~ Ò! t§~¡~¸¿¹È!Ñ9KÑ'LÐQRÒ&Rð 	0Ø/ó	0ÐRÐ&Rr"   c           	      ó^  — t        | j                  «      }t        d„ t        |d   «      D «       «      |d<   dj	                  t        |j                  «       «      D ��cg c]  \  }}|› d|› �‘Œ c}}«      }t        j                  |j                  d«      «      j                  «       S c c}}w )Nc              3   ó<   K  — | ]  \  }}|t        |«      f–— Œ y ­wr&   )rZ   )Ú.0ÚkÚvs      r    ú	<genexpr>z#CUDAOptions.hash.<locals>.<genexpr>‰   s   è ø€ Ò(h¹t¸qÀ!¨!¬Y°q«\Ô):Ñ(hùs   ‚r|   Ú_ú-r+   )
r†   Ú__dict__r‹   ÚsortedÚjoinrŒ   rU   rV   ÚencoderX   )r�   Ú	hash_dictÚnameÚvalÚkeys        r    ÚhashzCUDAOptions.hash‡   sŠ   € Ü˜Ÿ™Ó'ˆ	Ü#(Ñ(hÄvÈiÐXeÑNfÓGgÔ(hÓ#hˆ	�-Ñ Ø�h‰h¼À	ÇÁÓ@QÓ9R×S©I¨D°#˜4˜&  # šÓSÓTˆÜ�~‰~˜cŸj™j¨Ó1Ó2×<Ñ<Ó>Ð>ùó Ts   ÁB)
)"Ú__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Ø'-Ð Ó-Ø/IÐ  %¨¡*ÓIØ*.Ð! 4Ó.Ø€K�ÓØ€Eˆ4ÓØ€L�#ÓØ"Ð�tÓ"Ø€Dˆ#Óò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j4                  «       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        |j	                  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”   s@   € ØˆÜ—‘˜W dÓ+ˆÙÜÐGÈÀyÐQÓRÐRÜ�5—;‘;˜q“>Ó"Ð"r"   r   c                 óB   — | j                  |j                  «      }d|› �S )Núcuda:)r²   rG   )r�   rJ   r[   s      r    Úget_target_namezCUDABackend.get_target_name›   s#   € Ø×%Ñ% g§l¡lÓ3ˆ
Ø�z�lÐ#Ð#r"   Nc                 ó2   •— t         ‰| �  |«       d| _        y )NÚcubin)ÚsuperÚ__init__Ú
binary_ext)r�   r   Ú	__class__s     €r    r¹   zCUDABackend.__init__Ÿ   s   ø€ Ü‰Ñ˜Ô Ø!ˆ�r"   c                 óh  — dt         j                  j                  xs d| j                  j                  › �i}|j                  t        j                  j                  «       D �ci c]  }||v sŒ||   €Œ|||   “Œ c}«       t        | j                  |d   «      «      }d|vrFt        t        j                  «      }|dk\  r|j                  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 c c}w )NrG   Úsmrt   é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�   ÚoptsÚargsr“   r[   rt   s         r    Úparse_optionszCUDABackend.parse_options£   s.  € ØœŸ™×3Ñ3ÒN¸¸D¿K¹K×<LÑ<LÐ;MÐ7NÐOˆØ�‰¬×)IÑ)I×)NÑ)NÓ)PÖu AÐTUÐY]ÒT]ÐaeÐfgÑahÑat�Q˜˜Q™‘ZÒuÔvÜ˜×)Ñ)¨$¨v©,Ó7Ó8ˆ
à!¨Ñ-Ü#&¤{×'GÑ'GÓ#HÐ Ø˜RÒØ$×(Ñ(¨Ô3Ü+0´Ð8LÓ1MÓ+NˆDÐ'Ñ(à.°dÑ:Ø˜RÒØ<J�Ð8Ñ9à TÑ)Ü',§~¡~×'GÑ'GˆDÐ#Ñ$à9CÀrÒ9I±ÈqˆÐ,Ñ-äÑ"˜TÑ"Ð"ùò% vs   Á#	D/Á-D/Á3D/c                 óœ   — |j                   |j                  |j                  |j                  d   |j                  d   |j                  d   fS )Nr   r9   é   )rg   rh   Úsharedrl   )r�   Úmetadatas     r    Úpack_metadatazCUDABackend.pack_metadata¹   sO   € à×ÑØ×ÑØ�O‰OØ×!Ñ! !Ñ$Ø×!Ñ! !Ñ$Ø×!Ñ! !Ñ$ð
ð 	
r"   c                 óÌ   — dd l mc mc m} t	        | j                  |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Ã   sV   € ß1Ó1Ü˜×)Ñ)¨'¯,©,Ó7Ó8ˆ
ð 0:¸RÒ/?ˆD×+Ò+ÀT×EdÑ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Ø1°9Ð=Ð=r"   c                 ó.   — t        j                  |«       y r&   )r	   Úload_dialects)r�   Úctxs     r    rÜ   zCUDABackend.load_dialectsÑ   s   € Ü×Ñ˜SÕ!r"   c                 óÄ  — t        j                  | j                  «      }|j                  «        t        j
                  j                  |«       t        j                  j                  |«       |dz  dk  rt        j                  j                  |«       t        j
                  j                  |«       t        j                  j                  |«       t        j                  j                  |«       t        j
                  j                  |«       t        j
                  j                  |«       t        j                  j                  |«       |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Ø˜Ñ˜aÒÜ�K‰K×@Ñ@ÀÔDÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ Ô#Ü�‰×)Ñ)¨"Ô-Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ü�‰×#Ñ# BÔ'Ø
�‰ˆsŒØˆ
r"   c                 ó  — |j                   �H| j                  dt        j                  | j                  «      j                  |j                   «      «       t        j                  «       }|j                  �<|j                  d   |_	        |j                  d   |_
        |j                  d   |_        t        j                  | j                  «      }|j                  «       }t        j                  j!                  |d|› �|j"                  d|j$                  «       t        j&                  j)                  |«       |dz  dk\  rt        j&                  j+                  |«       t        j                  j,                  j/                  ||«       t        j&                  j1                  |«       t        j&                  j3                  |«       t        j&                  j5                  |«       t        j&                  j1                  |«       t        j&                  j7                  ||d	k\  «       t        j                  j,                  j9                  |«       t        j                  j;                  |«       |dz  d
v �rFt        j&                  j=                  |«       t        j>                  jA                  |«       t        j                  jC                  |«       t        j>                  jA                  |«       t        j&                  jE                  |«       t        j                  jF                  jI                  ||jJ                  |«       t        j&                  jM                  ||jJ                  «       t        j&                  jO                  |«       t        j&                  jQ                  ||jJ                  |«       �nÓ|dz  dk\  �r«t        j&                  j=                  |«       t        j>                  jA                  |«       t        j                  jC                  |«       t        j&                  jS                  |«       t        j&                  jU                  |«       t        j                  j,                  jW                  |«       t        j&                  jM                  ||jJ                  «       t        j&                  jO                  |«       t        j&                  jY                  ||jJ                  «       t        j&                  jQ                  ||jJ                  |«       t        j&                  jE                  |«       t        j                  j,                  j[                  |«       nt        j                  jC                  |«       t        j>                  jA                  |«       t        j                  j;                  |«       t        j&                  j]                  |«       t        j&                  j7                  ||d	k\  «       t        j&                  j_                  |«       t        j                  j,                  ja                  |«       t        j&                  j1                  |«       t        j                  j,                  jc                  |«       t        j&                  je                  |«       t        j&                  jg                  |«       t        j                  j;                  |«       t        j>                  ji                  |«       |dz  dk\  rRt        j                  j,                  jk                  |«       t        j                  j,                  jm                  |«       t        j>                  jo                  |«       t        j>                  jA                  |«       |jq                  | «       |j                  |j                  |j                  f|d<   | js                  «       }||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Ñ'7¸Ñ':ˆLÔ$Ø'*×'7Ñ'7¸Ñ':ˆLÔ$Ø'*×'7Ñ'7¸Ñ':ˆLÔ$Ü�_‰_˜SŸ[™[Ó)ˆØ—‘Ó(ˆÜ�‰×*Ñ*¨2°°z°lÐ/CÀSÇ]Á]ÐTVÐX[×XdÑXdÔeä�‰×#Ñ# BÔ'Ø˜Ñ˜qÒ Ü�N‰N×)Ñ)¨"Ô-ä�‰×Ñ×,Ñ,¨R°Ô>Ü�‰×4Ñ4°RÔ8Ü�‰×3Ñ3°BÔ7Ü�‰×,Ñ,¨RÔ0Ü�‰×4Ñ4°RÔ8Ü�‰×0Ñ0°°ZÀ2Ñ5EÔFÜ�‰×Ñ×@Ñ@ÀÔDÜ�‰×&Ñ& rÔ*Ø˜Ñ˜vÒ%Ü�N‰N×0Ñ0°Ô4Ü�M‰M×+Ñ+¨BÔ/Ü�K‰K×'Ñ'¨Ô+Ü�M‰M×+Ñ+¨BÔ/Ü�N‰N×;Ñ;¸BÔ?Ü�M‰M× Ñ ×4Ñ4°R¸¿¹ÈÔVÜ�N‰N×/Ñ/°°C·N±NÔCÜ�N‰N×-Ñ-¨bÔ1Ü�N‰N×'Ñ'¨¨C¯N©N¸LÖIØ˜2Ñ Ó#Ü�N‰N×0Ñ0°Ô4Ü�M‰M×+Ñ+¨BÔ/Ü�K‰K×'Ñ'¨Ô+Ü�N‰N×8Ñ8¸Ô<Ü�N‰N×/Ñ/°Ô3Ü�M‰M×#Ñ#×;Ñ;¸BÔ?Ü�N‰N×/Ñ/°°C·N±NÔCÜ�N‰N×-Ñ-¨bÔ1Ü�N‰N×.Ñ.¨r°3·>±>ÔBÜ�N‰N×'Ñ'¨¨C¯N©N¸LÔIÜ�N‰N×;Ñ;¸BÔ?Ü�M‰M×#Ñ#×:Ñ:¸2Õ>ä�K‰K×'Ñ'¨Ô+Ü�‰×'Ñ'¨Ô+Ü�‰×&Ñ& rÔ*Ü�‰×#Ñ# BÔ'Ü�‰×0Ñ0°°ZÀ2Ñ5EÔFÜ�‰×.Ñ.¨rÔ2Ü�‰×Ñ×9Ñ9¸"Ô=Ü�‰×4Ñ4°RÔ8Ü�‰×Ñ×3Ñ3°BÔ7Ü�‰×2Ñ2°2Ô6Ü�‰×/Ñ/°Ô3Ü�‰×&Ñ& rÔ*Ü�‰×$Ñ$ RÔ(Ø˜Ñ˜qÒ Ü�M‰M×#Ñ#×4Ñ4°RÔ8Ü�M‰M×#Ñ#×7Ñ7¸Ô;Ü�‰×Ñ˜rÔ"Ü�‰×'Ñ'¨Ô+Ø
�‰ˆsŒØ$0×$<Ñ$<¸l×>VÑ>VÐXd×XpÑXpÐ#qˆ�Ñ Ø×5Ñ5Ó7ˆØ&5ˆÐ"Ñ#Øˆ
r"   c                 óæ  — |}t        j                  |j                  «      }|j                  «        t        j
                  j                  |«       t        j                  j                  |«       t        j                  j                  |«       t        j
                  j                  |«       t        j
                  j                  |«       |j                  |«       |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ˆÐ"Ñ#Øˆ
r"   c                 ó¦  — t        || j                  j                  «      }|}t        j                  |j
                  «      }|j                  «        t        j                  j                  j                  |«       t        j                  j                  |«       t        j                  j                  |«       t        j                  j                  |«       t        j                  j!                  |«       t        j                  j                  j#                  |«       t        j                  j%                  |«       t        j                  j                  j'                  |||«       t        j(                  j+                  |«       t        j(                  j-                  |«       t        j                  j                  j/                  |«       t        j                  j                  j1                  |«       t        j(                  j+                  |«       t        j(                  j-                  |«       t        j(                  j3                  |«       t4        j6                  j8                  st        j:                  j=                  |«       |j?                  |«       tA        jB                  «        tA        j
                  «       }t4        j6                  jD                  rtG        d«      ‚tA        jH                  ||«      }	tK        |«      }
tM        || j                  j                  «      }d}t        jN                  «        tA        jP                  |	||
|«       t        jR                  |	«       |jT                  r4|jT                  D ��cg c]  \  }}|‘Œ	 }}}tA        jV                  |	|«       tA        jX                  |	t@        jZ                  «       |j]                  d«      }|�||d<   |j]                  d«      |d<   |j]                  d«      |d<   |j]                  d	«      |d
<   |j]                  d«      |d<   t_        |	«      }~	~|S c c}}w )NzYAddress Sanitizer Error: Address sanitizer is currently only supported on the AMD backendúnvptx64-nvidia-cudaz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   Útripler�   r/   ÚpathsÚtotal_num_warpsÚrets                     r    Ú	make_llirzCUDABackend.make_llir?  s  € Ü2°7¸D¿K¹K×<LÑ<LÓMˆàˆä�_‰_˜SŸ[™[Ó)ˆØ
�‰Ôä�‰×Ñ×-Ñ-¨bÔ1Ü�‰×7Ñ7¸Ô;Ü�‰×/Ñ/°Ô3Ü�‰×$Ñ$ RÔ(Ü�‰×1Ñ1°"Ô5Ü�‰×Ñ×:Ñ:¸2Ô>Ü�‰×9Ñ9¸"Ô=Ü�‰×Ñ×+Ñ+¨B°
¸KÔHÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×Ñ×1Ñ1°"Ô5Ü�‰×Ñ×;Ñ;¸BÔ?Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ü× Ñ ×2Ò2Ü�M‰M×&Ñ& rÔ*Ø
�‰ˆsŒä×ÑÔÜ—,‘,“.ˆÜ×Ñ×(Ò(ÜØkómð mä—>‘> # wÓ/ˆÜ& zÓ2ˆÜ ¨¯©×)9Ñ)9Ó:ˆØ&ˆÜ×ÑÔÜ×Ñ˜x¨°°xÔ@Ü×#Ñ# HÔ-à×ÒØ.5×.AÑ.A×B™l˜t T’TÐBˆEÑBÜ×!Ñ! (¨EÔ2ä×Ñ˜X¤t×'7Ñ'7Ô8ð ×*Ñ*Ð+@ÓAˆØÐ&Ø$3ˆH�[Ñ!Ø ×-Ñ-¨lÓ;ˆ�ÑØ #× 0Ñ 0Ð1IÓ Jˆ�ÑØ*-×*:Ñ*:Ð;[Ó*\ˆÐ&Ñ'Ø+.×+;Ñ+;Ð<aÓ+bˆÐ'Ñ(Ü�(‹mˆØØØˆ
ùó# Cs   ÎQc           	      ó–  — 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   rD  rC  rP   rG  Únamess              r    Úmake_ptxzCUDABackend.make_ptx{  s$  € Ü2°3¸¿¹×8HÑ8HÓIˆà&ˆÜ& zÓ2ˆÜ  T§[¡[×%5Ñ%5Ó6ˆÜ×#Ñ# C¨°°xÀÀS×EYÑEYÐ[`Óaˆä—
‘
ÐFÈÓLˆÜ�5‹z˜QŠÐˆØ  ™8ˆ�Ñà$ b™Ð)¨¨;°r©>Ð*:Ð;ˆÜ�f‰fÐ*¨i¸°}Ð,EÀsÔRT×R^ÑR^Ô_ˆÜ�f‰fÐ'¨;°z°lÐ)CÀSÔPR×P\ÑP\Ô]ˆä�f‰fÐ+¨R°Ó5ˆÜ�<‰<×"Ò"ÜÐ4Ô5Ü�#ŒJØˆ
r"   c                 óÀ  — t        «       j                  }t        j                  ddd¬«      5 }t        j                  ddd¬«      5 }|j	                  |«       |j                  «        |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                  j                  d«      ng }|g|	¢|
¢d‘|¢|¢d|› �‘|j                  ‘d‘|‘}	 t!        j"                  |dd|¬«       t$        j                  j'                  |j                  «      rt%        j(                  |j                  «       t$        j                  j'                  |j                  «      rt%        j(                  |j                  «       t-        |d«      5 }|j/                  «       }d d d «       t$        j                  j'                  |«      rt%        j(                  |«       d d d «       d d d «       S # t         j*                  $ rï}t-        |j                  «      5 }|j/                  «       }d d d «       n# 1 sw Y   nxY wt$        j                  j'                  |j                  «      rt%        j(                  |j                  «       |j0                  dk(  rd}n2|j0                  dt2        j4                  z   k(  rd}nd|j0                  › �}t7        |› d› ddj9                  |«      › d�«      ‚d }~ww xY w# 1 sw Y   �ŒUxY w# 1 sw Y   �Œ&xY w# 1 sw Y   S xY w)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
   r4  r5  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‘  sô  € Ü“× Ñ ˆÜ×(Ñ(°¸CÈÔOð .	 ÐSWÜ×'Ñ'¨u¸3ÀvÔNð.	 ØRVØ�J‰J�sŒOØ�J‰JŒLØ—9‘9˜tÑ#ˆDäAF×ARÑAR×AdÒAd˜Ð&<Ñ=ÐkvÐjwˆIØ×-Ò-‘2°NÐ3CˆDÜ*¨:Ó6ˆDô 38·,±,×2PÒ2P˜=¨#Ñ.ÐVXˆKð ?B¿oºo §¡× 5Ñ 5°cÔ :ÐSUÐð ðØ!ðØ$(ðØ*.ðØ1<ðØ?PðØT_Ð`dÐ_eÐRfðØhl×hqÑhqðØswðàðˆIðLÜ—‘˜y°ÀÈdÕSÜ—7‘7—>‘> $§)¡)Ô,Ü—I‘I˜dŸi™iÔ(Ü—7‘7—>‘> $§)¡)Ô,Ü—I‘I˜dŸi™iÔ(ô$ �d˜DÓ!ð ! QØŸ™›�÷!ä�w‰w�~‰~˜dÔ#Ü—	‘	˜$”÷].	 ÷ .	 ð^ ˆøô+ ×0Ñ0ò LÜ˜$Ÿ)™)“_ð *¨Ø"Ÿ-™-›/�C÷*÷ *ñ *úä—7‘7—>‘> $§)¡)Ô,Ü—I‘I˜dŸi™iÔ(à—<‘< 3Ò&Ø?‘EØ—\‘\ S¬6¯>©>Ñ%9Ò9Ø4‘Eà=¸a¿l¹l¸^ÐL�Eä  E 7ð +5Ø58°Eð :3Ø36·8±8¸IÓ3FÐ2GÀrð"Kó Lð LûðLú÷"!ñ !ú÷W.	 ñ .	 ú÷ .	 ð^ ˆúsz   ­MÁCMÄ B)H4Ç	MÇL9Ç&<MÈ"MÈ4L6ÉL1ÉI6	É-	L1É6I?É;B6L1Ì1L6Ì6MÌ9MÌ>MÍM	ÍMÍMc                 ó  ‡ ‡‡— ‰ j                  ‰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<   y )Nc                 ó,   •— ‰j                  | |‰‰«      S r&   )rò   ©r#  rÏ   r[   rJ   r�   s     €€€r    ú<lambda>z(CUDABackend.add_stages.<locals>.<lambda>Ç  s   ø€ °4·>±>À#ÀxÐQXÐZdÓ3e€ r"   rå   c                 ó,   •— ‰j                  | |‰‰«      S r&   )r!  r|  s     €€€r    r}  z(CUDABackend.add_stages.<locals>.<lambda>È  s   ø€ °D·O±OÀCÈÐSZÐ\fÓ4g€ r"   Úttgirc                 ó,   •— ‰j                  | |‰‰«      S r&   )r$  r|  s     €€€r    r}  z(CUDABackend.add_stages.<locals>.<lambda>Ê  s   ø€ °D·N±NÀ3ÈÐRYÐ[eÓ4f€ r"   c                 ó,   •— ‰j                  | |‰‰«      S r&   )rH  r|  s     €€€r    r}  z(CUDABackend.add_stages.<locals>.<lambda>Ë  s   ø€ ¨t¯~©~¸cÀ8ÈWÐV`Ó/a€ r"   Úllirc                 óT   •— ‰j                  | |‰‰j                  j                  «      S r&   )rS  r   rG   ©r#  rÏ   rJ   r�   s     €€r    r}  z(CUDABackend.add_stages.<locals>.<lambda>Ì  s#   ø€ ¨d¯m©m¸CÀÈ7ÐTX×T_ÑT_×TdÑTdÓ.e€ r"   Úptxc                 óT   •— ‰j                  | |‰‰j                  j                  «      S r&   )ry  r   rG   r„  s     €€r    r}  z(CUDABackend.add_stages.<locals>.<lambda>Í  s#   ø€ °·±ÀÀXÈwÐX\×XcÑXc×XhÑXhÓ0i€ r"   r·   )r²   rG   r   ÚTRITONÚGLUON)r�   ÚstagesrJ   rÇ   r[   s   ` ` @r    Ú
add_stageszCUDABackend.add_stagesÄ  se   ú€ Ø×%Ñ% g§l¡lÓ3ˆ
Ø”x—‘Ò&ÝeˆF�6‰NÝgˆF�7ŠOØœŸ™Ò'ÝfˆF�7‰OÝaˆˆv‰Üeˆˆu‰ÜiˆˆwŠr"   c                 óL   — t        «       }|› d| j                  j                  › �S )Nr—   )r3   r   rG   )r�   r2   s     r    r    zCUDABackend.hashÏ  s&   € ä#Ó%ˆØ�˜!˜DŸK™K×,Ñ,Ð-Ð.Ð.r"   )r¡   r¢   r£   Ústaticmethodr   rª   r²   r?   rµ   r¹   r   rË   rÐ   rØ   r   r   rÚ   rÜ   rò   r!  r$  rH  rS  ry  rŠ  Ú	functoolsÚ	lru_cacher    Ú__classcell__)r»   s   @r    r§   r§   Ž   sÖ   ø„ àð( 	ò (ó ð(ò#ð$¨#ó $ð"˜yð "¨Tõ "ð# Só #ò,
òð>  S¨* _Ñ 5ó >ò"ð ñó ðð  ñHó ðHòTò:òxò,1òf	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¬   ra  rk  rf  r-   Úpathlibr   r$   Ú
NvidiaToolr)   rŽ  r3   r#   rF   rK   rQ   rZ   rb   re   r§   r(   r"   r    ú<module>r™     s-  ðß EÑ Eß 8Ó 8Ý Ý ,å !Û ß -Ó -Ý Û Û 	Û Û Û 	Û Ý ð#˜ó #ð�5×#Ñ#ó ð €×ÑÓñó ðð €×ÑÓði Sò ió ðið$°ó ð €×ÑÓð
 ò 
ó ð
ð €×Ñ�TÓñ4ó ð4ð
&¨ó &ñ �$Ô÷'?ð '?ó ð'?ôTD/�+õ D/r"   