Ë
    Úéñi%\  ã                   ó  — 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
jB                  fd„Z" ejF                  «       dde fd„«       Z$ ejF                  «       de fd„«       Z%de fd„Z& ejF                  «       de fd„«       Z' ejF                  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   é   )ÚscalarÚprimitive_bitwidth)Úlhs_typeÚrhs_typeÚlhs_bitwidthÚrhs_bitwidths       úa/var/www/pod-logistic/pod-ai/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#   Úarchr   c                 ót   — | dk\  rt         j                  j                  S t         j                  j                  S )Néd   )r
   r	   Úptxas_blackwellÚptxas)r&   s    r!   Ú	get_ptxasr+   "   s'   € Ø+/°3ª;Œ5�<‰<×'Ñ'ÐN¼E¿L¹L×<NÑ<NÐNr#   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)r&   Úmock_verÚversions      r!   Úget_ptxas_versionr5   &   sL   € ä�|‰|×,Ñ,€HØÐØˆÜ×%Ñ%¤y°£×';Ñ';¸[Ð&IÓJ×QÑQÐRYÓZ€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 |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.
    ú.é   é   éP   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Ó4�L€Eˆ5Ø�‚{Ø�1Š9Ø˜‘:Ðà˜‘: ‘>Ð!Ø�‚{Ø�E‰zÐØ�‚{Ø�E‰zÐà�‚{ØˆØ˜5 2™:¨Ñ+Ñ+¨eÑ3Ð3ä
ÐXÐ[gÑgÓ
hÐhr#   c                 ób   — | j                   }|€ t        |«      j                  }t        |«      }|S ©N)Úptx_versionr+   r4   rJ   )Úoptionsr&   rM   rF   s       r!   Úget_ptx_version_from_optionsrO   G   s2   € Ø×%Ñ%€KØÐÜ  “×.Ñ.ˆÜ% lÓ3ˆØÐr#   c                 ó@   — t        | |«      }t        d|«      }d|› �}|S )NéV   z+ptx)rO   Úmin)rN   r&   rM   Úllvm_ptx_versionÚfeaturess        r!   Úget_featuresrU   O   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)r1   Úfs     r!   Ú	file_hashr^   ]   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 )Nr@   ÚaÚ Úsm_© )r_   Úsuffixs     r!   Úsm_arch_from_capabilityrf   c   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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%y	)'Ú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_dtypesrd   Ú!deprecated_fp8_dot_operand_dtypesÚtf32Údefault_dot_input_precision)r{   Útf32x3ÚieeeÚbf16x3Úbf16x6Úallowed_dot_input_precisionsÚmax_num_imprecise_acc_defaultÚextern_libsÚdebugÚcudaÚbackend_nameÚsanitize_overflowr&   rb   Úinstrumentation_modec                 óÆ  — 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   r   znum_warps must be a power of 2)r   Ú__file__Úparentrƒ   ÚdictÚgetr
   r	   Úlibdevice_pathrB   ÚobjectÚ__setattr__ÚtupleÚitemsrk   )Ú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rL   )r^   )Ú.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”   rY   rZ   Úencoder\   )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__rk   r$   Ú__annotations__rl   rn   ro   rp   r   rM   r
   r	   Úptxas_optionsrq   rB   rr   rs   Úboolrt   ru   rv   ry   r   rz   r|   r�   r‚   rƒ   rŽ   r„   r†   r‡   r&   rˆ   r—   r¨   rd   r#   r!   ri   ri   i   s   … à€IˆsÓØ€HˆcÓØ€J�ÓØ€IˆsÓð "€GˆX�c‰]Ó!Ø€K�ÓØ!&§¡×!;Ñ!;€K�˜#‘Ó;Ø!%€K�˜#‘Ó%Ø!Ð�dÓ!Ø#Ð˜Ó#Ø$)Ð˜TÓ)Ø€J�ÓØ'<Ð˜% ™*Ó<Ø46Ð% u¨S¡zÓ6Ø'-Ð Ó-Ø/]Ð  %¨¡*Ó]Ø*.Ð! 4Ó.Ø€K�ÓØ€Eˆ4ÓØ€L�#ÓØ"Ð�tÓ"Ø€Dˆ#ÓØ "Ð˜#Ó"ò0ó?r#   ri   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j6                  «       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        |j	                  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�   s@   € ØˆÜ—‘˜W dÓ+ˆÙÜÐGÈÀyÐQÓRÐRÜ�5—;‘;˜q“>Ó"Ð"r#   r   c                 óB   — | j                  |j                  «      }d|› �S )Núcuda:)r»   r&   )r•   rN   r_   s      r!   Úget_target_namezCUDABackend.get_target_name¤   s#   € Ø×%Ñ% g§l¡lÓ3ˆ
Ø�z�lÐ#Ð#r#   c                 ó2   •— t         ‰| �  |«       d| _        y )NÚcubin)ÚsuperÚ__init__Ú
binary_ext)r•   r   Ú	__class__s     €r!   rÂ   zCUDABackend.__init__¨   s   ø€ Ü‰Ñ˜Ô Ø!ˆ�r#   c                 óÜ  — d|v r|d   dk(  rd|d<   dt         j                  j                  xs d| j                  j                  › �i}|j                  t        j                  j                  «       D �ci c]  }||v sŒ||   €Œ|||   “Œ c}«       t        | j                  |d   «      «      }|j                  dd«      dkD  r|d	k  rt        d
|› 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 )Nrˆ   ÚconsanTr„   r&   Úsmrl   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.ry   éY   Úfp8e4nvrz   )rx   rs   i   @r   r‚   rd   )r
   ÚruntimeÚoverride_archr   r&   Úupdateri   Ú__dataclass_fields__Úkeysr$   r»   r�   r·   Úsetry   Úaddr“   r¡   ÚlanguageÚdefault_fp_fusion)r•   ÚoptsÚargsr›   r_   ry   s         r!   Úparse_optionszCUDABackend.parse_options¬   s�  € à! TÑ)¨dÐ3IÑ.JÈhÒ.VØ ˆD�‰MàœŸ™×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ˆ
à�8‰8�J Ó" QÒ&¨:¸ª?Üð !6Ø6@°\ð BNðOó Qð Qð "¨Ñ-Ü#&¤{×'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   Á4	E)Á>E)ÂE)c                 óH   — |j                   |j                  |j                  fS rL   )rk   rl   Úshared)r•   Úmetadatas     r!   Úpack_metadatazCUDABackend.pack_metadataË   s%   € à×ÑØ×ÑØ�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   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Ò   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                  |«       t        j                  r t        j                  j                  |«       y y rL   )r	   Úload_dialectsr°   Úinstrumentation)r•   Úctxs     r!   rå   zCUDABackend.load_dialectsà   s2   € Ü×Ñ˜SÔ!Ü×&Ò&Ü×'Ñ'×5Ñ5°cÕ:ð '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!                  | 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Ø˜Ñ˜aÒÜ�K‰K×@Ñ@ÀÔDÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ Ô#Ü�‰×)Ñ)¨"Ô-Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ü�‰×#Ñ# BÔ'Ø
�‰ˆs�KÔ Øˆ
r#   c                 ó  — |j                   �H| j                  dt        j                  | j                  «      j                  |j                   «      «       t        j                  | j                  «      }|j                  «       }|dz  dk\  }t        j                  j                  |d|› �|j                  d|j                  «       t        j                  j                  |«       t        j                  j                  ||«       t         j                  j"                  j%                  |«       t        j                  j'                  |«       t        j                  j)                  |«       t        j                  j+                  |«       t        j                  j'                  |«       t        j                  j-                  ||dk\  «       t         j                  j"                  j/                  |«       t        j                  j1                  |«       |dz  dv �rFt        j                  j3                  |«       t        j4                  j7                  |«       t        j                  j9                  |«       t        j4                  j7                  |«       t        j                  j;                  |«       t         j                  j<                  j?                  ||j@                  |«       t        j                  jC                  ||j@                  «       t        j                  jE                  |«       t        j                  jG                  ||j@                  |«       �n|dz  dk\  �rët        j                  j3                  |«       t        j4                  j7                  |«       t        j                  j9                  |«       t        j                  jI                  |«       t        j                  jK                  |d«       t         j                  j"                  jM                  |«       t        j                  jC                  ||j@                  «       t        j                  jE                  |«       t        j                  jO                  ||j@                  «       t        j                  jG                  ||j@                  |«       t        j                  jQ                  |«       t        j                  j;                  |«       t        j                  jK                  |d	«       t         j                  j"                  jS                  |«       nt        j                  j9                  |«       t        j4                  j7                  |«       t        j                  j1                  |«       t        j                  jU                  |«       t        j                  j-                  ||dk\  «       t        j                  jW                  |«       t         j                  j"                  jY                  |«       |dz  d
k\  r)t         j                  j"                  j[                  |«       t        j                  j'                  |«       t         j                  j"                  j]                  |«       t        j                  j_                  |«       t        j                  ja                  |«       t        j                  j1                  |«       t        j4                  jc                  |«       t         j                  j"                  je                  ||«       t         j                  j"                  jg                  |«       t        j4                  ji                  |«       t        j4                  jk                  |«       t        j4                  j7                  |«       |jm                  | d«       | jo                  «       |d<   | S )Nzttg.maxnregr=   r   r½   r   r:   )r   ré   FTré   Ú
make_ttgirÚtensordesc_meta)8rp   Úset_attrr   Úbuilderrì   Úget_int32_attrrë   rí   r   rð   Úadd_convert_to_ttgpuirrk   rl   Ú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_warpspecrn   Ú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°°z°lÐ/CÀSÇ]Á]ÐTVÐX[×XdÑXdÔeä�‰×#Ñ# BÔ'Ü�‰×%Ñ% b¨'Ô2ä�‰×Ñ×,Ñ,¨RÔ0Ü�‰×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×/Ñ/°°EÔ:Ü�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×7Ñ7¸Ô;Ü�N‰N×;Ñ;¸BÔ?ä�N‰N×/Ñ/°°DÔ9Ü�M‰M×#Ñ#×:Ñ:¸2Õ>ä�K‰K×'Ñ'¨Ô+Ü�‰×'Ñ'¨Ô+Ü�‰×&Ñ& rÔ*Ü�‰×#Ñ# BÔ'Ü�‰×0Ñ0°°ZÀ2Ñ5EÔFÜ�‰×.Ñ.¨rÔ2Ü�‰×Ñ×9Ñ9¸"Ô=Ø˜Ñ˜qÒ Ü�M‰M×#Ñ#×4Ñ4°RÔ8Ü�‰×4Ñ4°RÔ8Ü�‰×Ñ×3Ñ3°BÔ7Ü�‰×2Ñ2°2Ô6Ü�‰×/Ñ/°Ô3Ü�‰×&Ñ& rÔ*Ü�‰×$Ñ$ RÔ(Ü�‰×Ñ×3Ñ3°B¸
ÔCÜ�‰×Ñ×-Ñ-¨bÔ1Ü�‰×Ñ˜rÔ"Ü�‰×Ñ˜bÔ!Ü�‰×'Ñ'¨Ô+à
�‰ˆs�LÔ!Ø&)×&AÑ&AÓ&CˆÐ"Ñ#Øˆ
r#   c                 óô  — |}t        j                  |j                  «      }|j                  «        t        j
                  j                  |«       t        j
                  j                  |«       t        j
                  j                  |«       t        j                  j                  j                  |«       t        j
                  j                  |«       t        j                  j                  |«       t        j                  j!                  |«       t        j
                  j                  |«       t        j"                  j%                  |«       |j'                  |d«       |j)                  «       |d<   |S )NÚgluon_to_ttgirrÿ   )r   rë   rì   rí   r   Úgluonrï   Úadd_infer_coalesced_encodingsÚadd_resolve_auto_encodingsr	   r  r   ró   rî   r&  rð   r  r  r  rù   r'  )r•   ÚsrcrØ   rN   r_   rú   rü   s          r!   r+  zCUDABackend.gluon_to_ttgir@  só   € ØˆÜ�_‰_˜SŸ[™[Ó)ˆØ
�‰Ôä�‰× Ñ  Ô$Ü�‰×2Ñ2°2Ô6Ü�‰×/Ñ/°Ô3Ü�‰×Ñ×0Ñ0°Ô4Ü�‰×&Ñ& rÔ*Ü�‰×Ñ˜rÔ"Ü�‰×&Ñ& rÔ*Ü�‰×&Ñ& rÔ*Ü�‰×7Ñ7¸Ô;à
�‰ˆsÐ$Ô%Ø&)×&AÑ&AÓ&CˆÐ"Ñ#Øˆ
r#   c                 ó²  — t        || 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"                  j%                  |«       t        j                  j"                  j'                  |«       t(        j*                  j,                  dk(  rt        j                  j/                  |«       t        j                  j1                  |«       t        j                  j"                  j3                  ||«       t4        j6                  r+t4        j6                  j9                  d||j
                  «       t        j                  j                  j;                  |||«       t        j<                  j?                  |«       t        j<                  jA                  |«       t        j                  j"                  jC                  |«       t        j                  j"                  jE                  |«       t        j<                  j?                  |«       t        j<                  jA                  |«       t        j<                  jG                  |«       t        j                  jI                  |«       t(        j*                  jJ                  s9t(        j*                  jL                  st        jN                  jQ                  |«       t4        j6                  r+t4        j6                  j9                  d||j
                  «       |jS                  |d«       t(        j*                  jL                  rÚt(        j*                  jJ                  s`t        j                  |j
                  «      }|j                  «        t        jN                  jQ                  |«       |jS                  |d«       t        j                  |j
                  «      }|j                  «        t        jN                  jU                  |«       |jS                  |d«       tW        jX                  «        tW        j
                  «       }t(        j*                  jZ                  rt]        d«      ‚tW        j^                  ||«      }	ta        |«      }
tc        || j                  j                  «      }d}t        jd                  «        tW        jf                  |	||
|«       |jh                  rt        jj                  |	«       |jl                  rIt        jn                  |	«      r4|jl                  D ��cg c]  \  }}|‘Œ	 }}}tW        jp                  |	|«       tW        jr                  |	tV        jt                  «       |jw                  d	«      }|�||d
<   |jw                  d«      |d<   |jw                  d«      |d<   |jw                  d«      |d<   |jw                  d«      |d<   |jw                  d«      xs d|d<   |jw                  d«      xs d|d<   ty        |	«      }~	~|S c c}}w )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-cudazttg.total-num-warpsrk   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_cfr,  rï   r	   Úadd_allocate_shared_memory_nvr  Úadd_allocate_tensor_memoryÚadd_check_matmul_two_ctar
   Úcompilationrˆ   Ú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_llvmÚdisable_line_infoÚ"dump_ir_extract_di_local_variablesÚllvmirÚadd_di_scoperù   Úadd_di_local_variabler   Úinit_targetsÚenable_asanrE   Ú	to_modulerf   rU   Úset_short_ptrÚattach_datalayoutrt   Úset_nvvm_reflect_ftzrƒ   Úhas_extern_depsÚlink_extern_libsÚoptimize_moduleÚOPTIMIZE_O3Úget_int_attrrB   )r•   r/  rØ   rN   r_   rM   rú   rü   rì   Úllvm_modÚprocrT   Útripler¥   r1   ÚpathsÚtotal_num_warpsÚrets                     r!   r3  zCUDABackend.make_llirS  sÉ  € Ü2°7¸D¿K¹K×<LÑ<LÓMˆàˆä�_‰_˜SŸ[™[Ó)ˆØ
�‰Ôä�‰×7Ñ7¸Ô;Ü�‰×/Ñ/°Ô3Ü�‰×$Ñ$ RÔ(Ü�‰× Ñ  Ô$Ü�‰×Ñ×;Ñ;¸BÀ
ÈKÔXÜ�‰×Ñ×:Ñ:¸2Ô>Ü�‰×Ñ×8Ñ8¸Ô<Ü×Ñ×1Ñ1°XÒ=ä�N‰N×4Ñ4°RÔ8Ü�‰×9Ñ9¸"Ô=Ü�‰×Ñ×9Ñ9¸"¸jÔIä×&Ò&Ü×'Ñ'×-Ñ-Ð.AÀ2ÀsÇ{Á{ÔSÜ�‰×Ñ×+Ñ+¨B°
¸KÔHÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×Ñ×1Ñ1°"Ô5Ü�‰×Ñ×;Ñ;¸BÔ?Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ü�‰×'Ñ'¨Ô+ä× Ñ ×2Ò2¼5×;LÑ;L×;oÒ;oÜ�M‰M×&Ñ& rÔ*ä×&Ò&Ü×'Ñ'×-Ñ-Ð.>ÀÀCÇKÁKÔPà
�‰ˆs�KÔ ä×Ñ×?Ò?ä×$Ñ$×6Ò6Ü—_‘_ S§[¡[Ó1�Ø—‘Ô!Ü—‘×*Ñ*¨2Ô.Ø—‘�sÐ9Ô:ô —‘ §¡Ó-ˆBØ�O‰OÔÜ�M‰M×/Ñ/°Ô3Ø�F‰F�3ÐFÔGô 	×ÑÔÜ—,‘,“.ˆÜ×Ñ×(Ò(ÜØkómð mä—>‘> # wÓ/ˆÜ& zÓ2ˆÜ ¨¯©×)9Ñ)9Ó:ˆØ&ˆÜ×ÑÔÜ×Ñ˜x¨°°xÔ@Ø×%Ò%Ü×'Ñ'¨Ô1à×Ò¤6×#9Ñ#9¸(Ô#CØ.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ˆÐ'Ñ(Ø+.×+;Ñ+;Ð<]Ó+^Ò+cÐbcˆÐ'Ñ(Ø,/×,<Ñ,<Ð=cÓ,dÒ,iÐhiˆÐ(Ñ)Ü�(‹mˆØØØˆ
ùó' Cs   ×#[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 )Nr4  znvptx-mad-wide-optFz(.visible .entry ([a-zA-Z_][a-zA-Z0-9_]*)r   r   r¥   r=   r7   z\.version \d+\.\d+z	.version )Úflagsz\.target sm_\d+z.target sm_z,\s*debug|debug,\s*rb   z // -----// NVPTX Dump //----- //)rO   r   r&   rf   rU   r   Útranslate_to_asmrs   rµ   ÚfindallÚlenÚsubÚ	MULTILINEr
   r@  rJ  r	   Ú
dump_nvptxÚprint)r•   r/  rØ   rû   r_   rM   r[  rZ  rT   r`  r^  Únamess               r!   Úmake_ptxzCUDABackend.make_ptx³  s=  € Ü2°3¸¿¹×8HÑ8HÓIˆà&ˆÜ& zÓ2ˆÜ  T§[¡[×%5Ñ%5Ó6ˆØ%Ð&ˆÜ×#Ñ# C¨°°xÀÈ×H\ÑH\Ð^cÓdˆä—
‘
Ð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\Ô]ˆÜ× Ñ ×CÒCô —&‘&Ð/°°SÓ9ˆCÜ�<‰<×"Ò"ÜÐ4Ô5Ü�#ŒJØˆ
r#   c           
      ó$  — t        | j                  j                  «      j                  }t	        j
                  ddd¬«      5 }t	        j
                  ddd¬«      5 }|j                  |«       |j                  «        |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                   j#                  d«      ng }|g|	¢|
¢d‘|¢|¢d|› �‘|j                  ‘d‘|‘}	 t%        j&                  |dd|¬«       t        j                  j(                  r7t+        |j                  «      5 }t-        |j/                  «       «       d d d «       t0        j                  j3                  |j                  «      rt1        j4                  |j                  «       t0        j                  j3                  |j                  «      rt1        j4                  |j                  «       t+        |d«      5 }|j/                  «       }d d d «       t0        j                  j3                  |«      rt1        j4                  |«       d d d «       d d d «       S # 1 sw Y   �ŒxY w# t$        j6                  $ �r}t+        |j                  «      5 }|j/                  «       }d d d «       n# 1 sw Y   nxY wt0        j                  j3                  |j                  «      rt1        j4                  |j                  «       |j8                  dk(  rd}n2|j8                  dt:        j<                  z   k(  rd}nd|j8                  › �}|› d› ddj?                  |«      › d�}t-        d|› d|› d�«       tA        |«      ‚d }~ww xY w# 1 sw Y   �ŒwxY w# 1 sw Y   �ŒHxY w# 1 sw Y   S xY w) NFÚwz.ptx)ÚdeleteÚmodere   Úrz.logz.oz	-lineinfoz-suppress-debug-infoz-gz--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: ú
zC

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

zy
================================================================
please share the reproducer above with Triton project.
rW   )!r+   r   r&   r1   ÚtempfileÚNamedTemporaryFileÚwriteÚflushr¥   r
   r@  rI  r	   Údisable_ptxas_optrs   rf   rq   rD   r/   rù   Údump_ptxas_logrX   rg  r[   ÚosÚexistsÚremoveÚCalledProcessErrorÚ
returncodeÚsignalÚSIGSEGVr¢   r   )r•   r/  rØ   rû   r_   r*   ÚfsrcÚflogÚfbinÚ
debug_infoÚfmadr&   Údisable_optÚptx_extra_optionsÚ	ptxas_cmdÚlog_fileÚeÚlogÚerrorr]   rÀ   s                        r!   Ú
make_cubinzCUDABackend.make_cubinÍ  s›  € Ü˜$Ÿ+™+×*Ñ*Ó+×0Ñ0ˆÜ×(Ñ(°¸CÈÔOð G	 ÐSWÜ×'Ñ'¨u¸3ÀvÔNðG	 ØRVØ�J‰J�sŒOØ�J‰JŒLØ—9‘9˜tÑ#ˆDàˆJÜ× Ñ ×2Ò2à˜{Ð,BÐCÑC‘
Ü—‘×/Ò/à˜t˜fÑ$‘
ð ˜{˜mÑ+�
à×-Ò-‘2°NÐ3CˆDÜ*¨:Ó6ˆDô 38·,±,×2PÒ2P˜=¨#Ñ.ÐVXˆKð ?B¿oºo §¡× 5Ñ 5°cÔ :ÐSUÐð ðØ"ðØ%)ðØ+/ðØ2=ðØ@QðØU`ÐaeÐ`fÐSgðØim×irÑirðàðàðˆIð$(Ü—‘˜y°ÀÈdÕSÜ—<‘<×.Ò.Ü˜dŸi™i›ð /¨HÜ˜hŸm™m›oÔ.÷/ô —7‘7—>‘> $§)¡)Ô,Ü—I‘I˜dŸi™iÔ(Ü—7‘7—>‘> $§)¡)Ô,Ü—I‘I˜dŸi™iÔ(ô: �d˜DÓ!ð ! QØŸ™›�÷!ä�w‰w�~‰~˜dÔ#Ü—	‘	˜$”÷OG	 ÷ G	 ðP ˆ÷O/ñ /ûô ×0Ñ0ó (Ü˜$Ÿ)™)“_ð *¨Ø"Ÿ-™-›/�C÷*÷ *ñ *úä—7‘7—>‘> $§)¡)Ô,Ü—I‘I˜dŸi™iÔ(à—<‘< 3Ò&Ø?‘EØ—\‘\ S¬6¯>©>Ñ%9Ò9Ø4‘Eà=¸a¿l¹l¸^ÐL�Eà!˜7ð #-Ø-0¨Eð 2+Ø+.¯8©8°IÓ+>Ð*?¸rðC�ô ð ð €ð à€ð ðô ô ! Ó'Ð'ûð5(ú÷8!ñ !ú÷IG	 ñ G	 ú÷ G	 ðP ˆús”   ÁPÁDO8ÅAKÆ'KÇBKÉO8É%O+É6<O8Ê2PËKË	KËO(Ë%O#Ë:L	Ì	O#ÌLÌC
O#Ï#O(Ï(O8Ï+O5Ï0O8Ï8P	Ï=PÐP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<   t
        j                  j                  �$t
        j                  j                  ‰ |‰|‰«       y y )Nc                 ó,   •— ‰j                  | |‰‰«      S rL   )rê   ©r/  rØ   r_   rN   r•   s     €€€r!   ú<lambda>z(CUDABackend.add_stages.<locals>.<lambda>  s   ø€ °4·>±>À#ÀxÐQXÐZdÓ3e€ r#   rð   c                 ó,   •— ‰j                  | |‰‰«      S rL   )rþ   r“  s     €€€r!   r”  z(CUDABackend.add_stages.<locals>.<lambda>  s   ø€ °D·O±OÀCÈÐSZÐ\fÓ4g€ r#   Úttgirc                 ó,   •— ‰j                  | |‰‰«      S rL   )r+  r“  s     €€€r!   r”  z(CUDABackend.add_stages.<locals>.<lambda>  s   ø€ °D×4GÑ4GÈÈXÐW^Ð`jÓ4k€ r#   c                 ó,   •— ‰j                  | |‰‰«      S rL   )r3  r“  s     €€€r!   r”  z(CUDABackend.add_stages.<locals>.<lambda>   s   ø€ ¨t¯~©~¸cÀ8ÈWÐV`Ó/a€ r#   Úllirc                 óT   •— ‰j                  | |‰‰j                  j                  «      S rL   )ri  r   r&   ©r/  rØ   rN   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 rL   )r�  r   r&   r›  s     €€r!   r”  z(CUDABackend.add_stages.<locals>.<lambda>"  s#   ø€ °·±ÀÀXÈwÐX\×XcÑXc×XhÑXhÓ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¡lÓ3ˆ
Ø”x—‘Ò&ÝeˆF�6‰NÝgˆF�7ŠOØœŸ™Ò'ÝkˆF�7‰OÝaˆˆv‰Üeˆˆu‰Üiˆˆw‰Ü�=‰=×3Ñ3Ð?Ü�M‰M×4Ñ4°T¸6À7ÈHÐV`Õað @r#   c                 óv   — t        | j                  j                  «      }|› d| j                  j                  › �S )NrŸ   )r5   r   r&   )r•   r4   s     r!   r¨   zCUDABackend.hash&  s2   € ä# D§K¡K×$4Ñ$4Ó5ˆØ�˜!˜DŸK™K×,Ñ,Ð-Ð.Ð.r#   )r©   rª   r«   ræ   Ústaticmethodr   r³   r»   rB   r¾   rÂ   r   rÕ   rÙ   rá   r   r   rã   rå   rê   rþ   r+  r3  ri  r�  r¢  Ú	functoolsÚ	lru_cacher¨   Ú__classcell__)rÄ   s   @r!   r°   r°   –   sÝ   ø„ Ø€Oàð( 	ò (ó ð(ò#ð$¨#ó $ð"˜yð "¨Tõ "ð# Só #ò>
òð>  S¨* _Ñ 5ó >ò;ð
 ñó ðð  ñGó ðGòRò&^ò@ò4JòX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   rY   rµ   rw  r‚  r}  r/   Úpathlibr   r%   r$   Ú
NvidiaToolr+   r¦  r5   rJ   rO   rU   r^   rf   ri   r°   rd   r#   r!   ú<module>r±     s>  ðß EÑ Eß 8Ó 8Ý Ý ,å !Û ß -Ó -Ý Û Û 	Û Û Û 	Û Ý ð#˜ó #ðO�Cð O˜E×,Ñ,ó Oð €×ÑÓñ˜Cò ó ðð €×ÑÓði Sò ió ðið.°ó ð €×ÑÓð
 ò 
ó ð
ð €×Ñ�TÓñ4ó ð4ð
&¨ó &ñ �$Ô÷)?ð )?ó ð)?ôXS/�+õ S/r#   