Ë
    0²j1]  ã                   ó  — 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       ú^/var/www/html/ia-enhanced/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 )Nr@   z+ptx)rO   Úmin)rN   r&   rM   Úllvm_ptx_versionÚfeaturess        r!   Úget_featuresrT   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_capabilityre   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%e&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_dtypesrc   Ú!deprecated_fp8_dot_operand_dtypesÚtf32Údefault_dot_input_precision)rz   Útf32x3ÚieeeÚbf16x3Úbf16x6Úallowed_dot_input_precisionsÚmax_num_imprecise_acc_defaultÚextern_libsÚdebugÚcudaÚbackend_nameÚsanitize_overflowr&   ra   Úinstrumentation_modec                 óÆ  — t        t        «      j                  dz  }| j                  €i nt	        | j                  «      }|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Úitemsrj   )Ú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“   rX   rY   Ú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)
c                 ó   — d| j                   v S )NÚiisan)r‡   )r”   s    r!   Úenable_iisanzCUDAOptions.enable_iisan•   s   € à˜$×3Ñ3Ð3Ð3r#   )(Ú__name__Ú
__module__Ú__qualname__rj   r$   Ú__annotations__rk   rm   rn   ro   r   rM   r
   r	   Úptxas_optionsrp   rB   rq   rr   Úboolrs   rt   ru   rx   r   ry   r{   r€   r�   r‚   r�   rƒ   r…   r†   r&   r‡   r–   r§   Úpropertyrª   rc   r#   r!   rh   rh   i   s4  … à€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ò?ð ñ4ó ñ4r#   rh   c                   óô   ‡ — e Zd ZdZedefd„«       Zd„ Zdefd„Z	deddfˆ fd„Z
defd„Zd	„ Zd
„ Zdeeef   fd„Zd„ Zed„ «       Zed„ «       Zd„ Zd„ Zd„ Zd„ Zd„ Z e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                 óø  ‡— t        ˆfd„dD «       «      r
d‰d<   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 )Nc              3   óF   •K  — | ]  }|‰j                  d d«      v –— Œ y­w)r‡   ra   N)rŽ   )r™   ÚmodeÚoptss     €r!   rœ   z,CUDABackend.parse_options.<locals>.<genexpr>²   s"   øè ø€ Ò\Àˆt�t—x‘xÐ 6¸Ó;Ô;Ñ\ùs   ƒ!)Úconsanr©   Trƒ   Fr†   r&   Úsmrk   r   r@   zBnum_ctas > 1 requires NVIDIA SM90+ (Hopper). Current target is sm_zM. This configuration will fail. Please set num_ctas=1 or target an SM90+ GPU.rx   éY   Úfp8e4nvry   )rw   rr   i   @r   r�   rc   )Úanyr
   ÚruntimeÚoverride_archr   r&   Úupdaterh   Ú__dataclass_fields__Úkeysr$   r¾   rŽ   rº   Úsetrx   Úaddr’   r    ÚlanguageÚdefault_fp_fusion)r”   rË   Úargsrš   r^   rx   s    `    r!   Úparse_optionszCUDABackend.parse_options°   s“  ø€ äÓ\ÐH[Ô\Ô\Ø ˆD�‰MØ(-ˆDÐ$Ñ%àœŸ™×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   Â	E7ÂE7ÂE7c                 óH   — |j                   |j                  |j                  fS rL   )rj   rk   Ú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)8ro   Úset_attrr   Úbuilderrò   Úget_int32_attrrñ   ró   r   rö   Úadd_convert_to_ttgpuirrj   rk   ÚttgpuirÚadd_coalesceÚadd_f32_dot_tcr	   Ú	ttnvgpuirÚadd_plan_ctaÚadd_remove_layout_conversionsÚadd_optimize_thread_localityÚadd_accelerate_matmulÚadd_optimize_dot_operandsÚ add_optimize_descriptor_encodingÚadd_loop_aware_cseÚadd_fuse_nested_loopsrô   rù   Úadd_triton_licmÚ add_combine_tensor_select_and_ifÚhopperÚadd_hopper_warpspecrm   Úadd_assign_latenciesÚadd_schedule_loopsÚadd_pipelineÚadd_optimize_accumulator_initÚadd_hoist_tmem_allocÚadd_promote_lhs_to_tmemÚadd_warp_specializeÚadd_optimize_partition_warpsÚadd_remove_tmem_tokensÚadd_prefetchÚadd_coalesce_async_copyÚadd_optimize_tmem_layoutsÚadd_tma_loweringÚadd_interleave_tmemÚadd_reduce_data_duplicationÚadd_reorder_instructionsrý   Úadd_fence_insertionÚadd_lower_mmaÚadd_sccprü   rÿ   Úget_tensordesc_metadata)r   rÞ   r  r^   r  Údump_enabledÚemuTF32s          r!   r  zCUDABackend.make_ttgirû   s  € ð �;‰;Ð"Ø�L‰L˜¬¯
©
°3·;±;Ó(?×(NÑ(NÈsÏ{É{Ó([Ô\ä�_‰_˜SŸ[™[Ó)ˆØ—‘Ó(ˆØ Ñ# qÑ(ˆÜ�‰×*Ñ*¨2°°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!   r1  zCUDABackend.gluon_to_ttgirE  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'                  |«       d|j(                  v rt        j                  j+                  |«       t        j                  j-                  |«       t        j                  j"                  j/                  ||«       t0        j2                  r+t0        j2                  j5                  d||j
                  «       t        j                  j                  j7                  |||«       t        j8                  j;                  |«       t        j8                  j=                  |«       t        j                  j"                  j?                  |«       t        j                  j"                  jA                  |«       t        j8                  j;                  |«       t        j8                  j=                  |«       t        j8                  jC                  |«       t        j                  jE                  |«       tF        jH                  jJ                  s9tF        jH                  jL                  st        jN                  jQ                  |«       t0        j2                  r+t0        j2                  j5                  d||j
                  «       |jS                  |d«       tF        jH                  jL                  rÚtF        jH                  jJ                  s`t        j                  |j
                  «      }|j                  «        t        jN                  jQ                  |«       |jS                  |d«       t        j                  |j
                  «      }|j                  «        t        jN                  jU                  |«       |jS                  |d«       tW        jX                  «        tW        j
                  «       }tF        jH                  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-warpsrj   z
ttg.sharedrÝ   zttg.tensor_memory_sizeÚ	tmem_sizezttg.global_scratch_memory_sizeÚglobal_scratch_sizez#ttg.global_scratch_memory_alignmentÚglobal_scratch_alignzttg.profile_scratch_memory_sizer   Úprofile_scratch_sizez$ttg.profile_scratch_memory_alignmentr   Úprofile_scratch_align)=rO   r   r&   r   rñ   rò   ró   r   r
  r  Úadd_allocate_warp_groupsÚconvertÚadd_scf_to_cfr2  rõ   r	   Úadd_allocate_shared_memory_nvr  Úadd_allocate_tensor_memoryÚadd_check_matmul_two_ctar‡   Úadd_concurrency_sanitizerÚ"add_allocate_global_scratch_memoryÚadd_proxy_fence_insertionr³   rì   ÚpatchÚadd_to_llvmirrô   rù   rü   Úadd_nvgpu_to_llvmÚadd_warp_specialize_to_llvmrý   Úadd_nvvm_to_llvmr
   ÚcompilationÚdisable_line_infoÚ"dump_ir_extract_di_local_variablesÚllvmirÚadd_di_scoperÿ   Úadd_di_local_variabler   Úinit_targetsÚenable_asanrE   Ú	to_modulere   rT   Úset_short_ptrÚattach_datalayoutrs   Úset_nvvm_reflect_ftzr‚   Úhas_extern_depsÚlink_extern_libsÚoptimize_moduleÚOPTIMIZE_O3Úget_int_attrrB   )r”   r5  rÞ   rN   r^   rM   r   r  rò   Úllvm_modÚprocrS   Útripler¤   r1   ÚpathsÚtotal_num_warpsÚrets                     r!   r9  zCUDABackend.make_llirX  sÃ  € Ü2°7¸D¿K¹K×<LÑ<LÓMˆàˆä�_‰_˜SŸ[™[Ó)ˆØ
�‰Ôä�‰×7Ñ7¸Ô;Ü�‰×/Ñ/°Ô3Ü�‰×$Ñ$ RÔ(Ü�‰× Ñ  Ô$Ü�‰×Ñ×;Ñ;¸BÀ
ÈKÔXÜ�‰×Ñ×:Ñ:¸2Ô>Ü�‰×Ñ×8Ñ8¸Ô<Ø�w×3Ñ3Ñ3ä�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 )Nr:  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*ra   z // -----// NVPTX Dump //----- //)rO   r   r&   re   rT   r   Útranslate_to_asmrr   r¸   ÚfindallÚlenÚsubÚ	MULTILINEr
   rN  rP  r	   Ú
dump_nvptxÚprint)r”   r5  rÞ   r  r^   rM   ra  r`  rS   rf  rd  Ú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 }d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)ÚdeleterÊ   rd   Úrz.logz.oz	-lineinfoz-suppress-debug-infoz-gz--fmad=falsez--opt-levelÚ0ú z--regAllocOptLevel=2z-vz--gpu-name=z-oT)ÚcheckÚ	close_fdsÚstderréÿ   z!Internal Triton PTX codegen erroré€   z`ptxas` raised SIGSEGVz`ptxas` failed with error code z
`ptxas` stderr:
z
Repro command: ú
zC

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

zy
================================================================
please share the reproducer above with Triton project.
rV   )!r+   r   r&   r1   ÚtempfileÚNamedTemporaryFileÚwriteÚflushr¤   r
   rN  rO  r	   Údisable_ptxas_optrr   re   rp   rD   r/   rÿ   Údump_ptxas_logrW   rm  rZ   ÚosÚexistsÚremoveÚCalledProcessErrorÚ
returncodeÚsignalÚSIGSEGVr¡   r   )r”   r5  rÞ   r  r^   r*   ÚfsrcÚflogÚfbinÚ
debug_infoÚfmadr&   Údisable_optÚptx_extra_optionsÚ	reg_allocÚ	ptxas_cmdÚlog_fileÚeÚlogÚerrorr\   rÃ   s                         r!   Ú
make_cubinzCUDABackend.make_cubinÒ  s«  € Ü˜$Ÿ+™+×*Ñ*Ó+×0Ñ0ˆÜ×(Ñ(°¸CÈÔOð J	 ÐSWÜ×'Ñ'¨u¸3ÀvÔNðJ	 Ø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Ðð 0Ð0ˆIð ðØ"ðØ%)ðØ+/ðØ2=ðØ@IðØL]ðØalÐmqÐlrÐ_sðà—	‘	ðàðà!%ðˆ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Ô#Ü—	‘	˜$”÷UJ	 ÷ J	 ðV ˆ÷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!ñ !ú÷OJ	 ñ J	 ú÷ J	 ðV ˆús”   ÁP
ÁDO=Å$AKÆ,K	ÇBKÉO=É*O0É;<O=Ê7P
Ë	KËKËO-Ë*O(Ë?L	Ì	O(ÌL"ÌC
O(Ï(O-Ï-O=Ï0O:Ï5O=Ï=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ð   ©r5  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   )r1  r™  s     €€€r!   rš  z(CUDABackend.add_stages.<locals>.<lambda>'  s   ø€ °D×4GÑ4GÈÈXÐW^Ð`jÓ4k€ r#   c                 ó,   •— ‰j                  | |‰‰«      S rL   )r9  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   )ro  r   r&   ©r5  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  r1  r9  ro  r–  r¨  Ú	functoolsÚ	lru_cacher§   Ú__classcell__)rÇ   s   @r!   r³   r³   š   sÞ   ø„ Ø€Oàð( 	ò (ó ð(ò#ð$¨#ó $ð"˜yð "¨Tõ "ð# Só #ò@
òð>  S¨* _Ñ 5ó >ò;ð
 ñó ðð  ñGó ðGòRò&^ò@ò4Mò^bð €Y×ÑÓñ/ó ô/r#   r³   )r:   ),Útriton.backends.compilerr   r   r   Útriton._C.libtritonr   r   r   r	   Útritonr
   Útriton.runtime.errorsr   Údataclassesr   r«  Útypingr   r   r   r   Útypesr   rX   r¸   r|  r‡  r‚  r/   Úpathlibr   r%   r$   Ú
NvidiaToolr+   r¬  r5   rJ   rO   rT   r]   re   rh   r³   rc   r#   r!   ú<module>r·     s>  ðß EÑ Eß 8Ó 8Ý Ý ,å !Û ß -Ó -Ý Û Û 	Û Û Û 	Û Ý ð#˜ó #ðO�Cð O˜E×,Ñ,ó Oð €×ÑÓñ˜Cò ó ðð €×ÑÓði Sò ió ðið.°ó ð €×ÑÓð
 ò 
ó ð
ð €×Ñ�TÓñ4ó ð4ð
&¨ó &ñ �$Ô÷-4ð -4ó ð-4ô`W/�+õ W/r#   