ó
    �~i%\  ã                   ó  • S SK JrJrJr  S SKJrJrJrJr  S SK	J
r
  S SKJr  S SKJr  S SKrS SKJrJrJrJr  S SKJr  S SKrS SKrS SKrS SKrS SKrS SKrS S	KJr  S
\4S jrS\ S\
RB                  4S jr"\RF                  " 5       SS\ 4S jj5       r$\RF                  " 5       S\ 4S j5       r%S\ 4S jr&\RF                  " 5       S\ 4S j5       r'\RF                  " S5      S 5       r(S\ 4S jr)\" SS9 " S S5      5       r* " S S\5      r+g)é    )ÚBaseBackendÚ	GPUTargetÚLanguage)ÚirÚpassesÚllvmÚnvidia)Úknobs)Ú
PTXASError)Ú	dataclassN)ÚAnyÚDictÚTupleÚOptional)Ú
ModuleType)ÚPathÚtargetc                 ó@   • S[         [        [        [        4   4S jnU$ )NÚreturnc                 ó‚   • U R                   R                  nUR                   R                  nX#:X  d   S5       eUS:X  a  gg)Nz%lhs and rhs bitwidth must be the sameé   )é   r   é    )r   r   é   )ÚscalarÚprimitive_bitwidth)Úlhs_typeÚrhs_typeÚlhs_bitwidthÚrhs_bitwidths       Ú\/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/backends/nvidia/compiler.pyÚcheck_dot_compatibilityÚ-min_dot_size.<locals>.check_dot_compatibility   s@   € Ø—‘×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                 óv   • U S:¼  a  [         R                  R                  $ [         R                  R                  $ )Néd   )r
   r	   Úptxas_blackwellÚptxas)r'   s    r!   Ú	get_ptxasr,   "   s'   € Ø+/°3«;Œ5�<‰<×'Ñ'ÐN¼E¿L¹L×<NÑ<NÐNr$   c                 ó¸   • [         R                  R                  nUb  U$ [        R                  " [        U 5      R                  S/5      R                  S5      nU$ )Nz	--versionúutf-8)r
   r	   Úmock_ptx_versionÚ
subprocessÚcheck_outputr,   ÚpathÚdecode)r'   Úmock_verÚversions      r!   Úget_ptxas_versionr6   &   sL   € ä�|‰|×,Ñ,€HØÑØˆÜ×%Ò%¤y°£×';Ñ';¸[Ð&IÓJ×QÑQÐRYÓZ€GØ€Nr$   c                 ó  • [        U [        5      (       d   e[        [        U R	                  S5      5      u  pUS:X  a  US:  a  SU-   $ SU-   S-
  $ US:X  a  SU-   $ US:X  a  S	U-   $ US
:¼  a  SnX1S
-
  S-  -   U-   $ [        SU -   5      e)zC
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_versionrK   /   s­   € ô
 �l¤C×(Ñ(Ð(Ð(Ü”s˜L×.Ñ.¨sÓ3Ó4�L€EØ�ƒ{Ø�1‹9Ø˜‘:Ðà˜‘: ‘>Ð!Ø�ƒ{Ø�E‰zÐØ�ƒ{Ø�E‰zÐà�ƒ{ØˆØ 2™:¨Ñ+Ñ+¨eÑ3Ð3ä
ÐXÐ[gÑgÓ
hÐhr$   c                 ód   • U R                   nUc   [        U5      R                  n[        U5      nU$ ©N)Úptx_versionr,   r5   rK   )Úoptionsr'   rN   rG   s       r!   Úget_ptx_version_from_optionsrP   G   s2   € Ø×%Ñ%€KØÑÜ  “×.Ñ.ˆÜ% lÓ3ˆØÐr$   c                 ó>   • [        X5      n[        SU5      nSU 3nU$ )NéV   z+ptx)rP   Úmin)rO   r'   rN   Úllvm_ptx_versionÚfeaturess        r!   Úget_featuresrV   O   s.   € ä.¨wÓ=€Kô ˜2˜{Ó+ÐØÐ&Ð'Ð(€HØ€Or$   c                 ó´   • [        U S5       n[        R                  " UR                  5       5      R	                  5       sS S S 5        $ ! , (       d  f       g = f)NÚrb)ÚopenÚhashlibÚsha256ÚreadÚ	hexdigest)r2   Úfs     r!   Ú	file_hashr_   ]   s5   € ä	ˆd�DÔ	˜QÜ�~Š~˜aŸf™f›hÓ'×1Ñ1Ó3÷ 
×	×	ús   �2A	Á	
AÚ
capabilityc                 ó$   • U S:¼  a  SOSnSU  U 3$ )NrA   ÚaÚ Úsm_© )r`   Úsuffixs     r!   Úsm_arch_from_capabilityrg   c   s!   € à "Ó$‰S¨"€FØ��˜V˜HÐ%Ð%r$   T)Úfrozenc                   ó°  • \ rS rSr% Sr\\S'   Sr\\S'   Sr\\S'   Sr	\\S	'   S
r
\\   \S'   S
r\\S'   \R                  R                  r\\   \S'   S
r\\   \S'   Sr\\S'   Sr\\S'   Sr\\S'   Sr\\S'   Sr\\   \S'   Sr\\   \S'   Sr\\S'   Sr\\   \S'   S
r\\S'   S
r\\S'   Sr \\S'   S r!\\S!'   Sr"\\S"'   S
r#\\S#'   S$r$\\S%'   S& r%S' r&Sr'g
)(ÚCUDAOptionséi   é   Ú	num_warpsr   Únum_ctasé   Ú
num_stagesr   Ú	warp_sizeNÚmaxnregrN   Úptx_optionsÚir_overrideTÚenable_fp_fusionÚenable_reflect_ftzFÚlaunch_cooperative_gridÚ
launch_pdl)Úfp8e5Úfp8e4b15Úsupported_fp8_dtypesre   Ú!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'   rc   Úinstrumentation_modec                 óà  • [        [        5      R                  S-  nU R                  c  0 O[	        U R                  5      nUR                  SS 5      (       d2  [        R                  R                  =(       d    [        US-  5      US'   [        R                  U S[        UR                  5       5      5        U R                  S:”  a   U R                  U R                  S-
  -  S:X  d   S5       eg )NÚlibÚ	libdevicezlibdevice.10.bcr…   r   r   znum_warps must be a power of 2)r   Ú__file__Úparentr…   ÚdictÚgetr
   r	   Úlibdevice_pathrC   ÚobjectÚ__setattr__ÚtupleÚitemsrm   )ÚselfÚdefault_libdirr…   s      r!   Ú__post_init__ÚCUDAOptions.__post_init__…   sÁ   € Üœh›×.Ñ.°Ñ6ˆØ ×,Ñ,Ñ4‘b¼$¸t×?OÑ?OÓ:PˆØ�‰˜{¨D×1Ñ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           	      ód  • [        U R                  5      n[        S [        US   5       5       5      US'   SR	                  [        UR                  5       5       VVs/ s H  u  p#U SU 3PM     snn5      n[        R                  " UR                  S5      5      R                  5       $ s  snnf )Nc              3   ó@   #   • U  H  u  pU[        U5      4v •  M     g 7frM   )r_   )Ú.0ÚkÚvs      r!   Ú	<genexpr>Ú#CUDAOptions.hash.<locals>.<genexpr>‘   s   é € Ð(hÒGg¹t¸q¨!¬Y°q«\Õ):ÒGgùs   ‚r…   Ú_Ú-r.   )
r�   Ú__dict__r•   ÚsortedÚjoinr–   rZ   r[   Úencoder]   )r—   Ú	hash_dictÚnameÚvalÚkeys        r!   ÚhashÚCUDAOptions.hash�   sŽ   € Ü˜Ÿ™Ó'ˆ	Ü#(Ñ(hÄvÈiÐXeÑNfÔGgÓ(hÓ#hˆ	�-Ñ Ø�h‰h¼À	ÇÁÓ@QÔ9RÔSÒ9R©I¨D˜4˜&  # ›Ñ9RÒSÓTˆÜ�~Š~˜cŸj™j¨Ó1Ó2×<Ñ<Ó>Ð>ùó Ts   ÁB,
)(Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__rm   r%   Ú__annotations__rn   rp   rq   rr   r   rN   r
   r	   Úptxas_optionsrs   rC   rt   ru   Úboolrv   rw   rx   r{   r   r|   r~   rƒ   r„   r…   r�   r†   rˆ   r‰   r'   rŠ   r™   r¬   Ú__static_attributes__re   r$   r!   rj   rj   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$   rj   c                   ó  ^ • \ rS rSrSr\S\4S j5       rS rS\	4S jr
S\SS4U 4S jjrS\4S	 jrS
 rS rS\\	\4   4S jrS r\S 5       r\S 5       rS rS rS rS rS r\R8                  " 5       S 5       rSrU =r$ )ÚCUDABackendé–   Nr   c                 ó    • U R                   S:H  $ )Nr‡   )Úbackend)r   s    r!   Úsupports_targetÚCUDABackend.supports_target™   s   € à�~‰~ Ñ'Ð'r$   c                 ó�   • Sn[         R                  " X!5      nU(       d  [        SU 35      e[        UR	                  S5      5      $ )Nz	^sm(\d+)$z(TRITON_OVERRIDE_ARCH must have the form r   )ÚreÚ	fullmatchÚ
ValueErrorr%   Úgroup)r—   r'   ÚpatternÚmatchs       r!   Ú_parse_archÚCUDABackend._parse_arch�   s>   € ØˆÜ—’˜WÓ+ˆÞÜÐGÈÀyÐQÓRÐRÜ�5—;‘;˜q“>Ó"Ð"r$   r   c                 óB   • U R                  UR                  5      nSU 3$ )Núcuda:)rÄ   r'   )r—   rO   r`   s      r!   Úget_target_nameÚCUDABackend.get_target_name¤   s#   € Ø×%Ñ% g§l¡lÓ3ˆ
Ø�z�lÐ#Ð#r$   c                 ó2   >• [         TU ]  U5        SU l        g )NÚcubin)ÚsuperÚ__init__Ú
binary_ext)r—   r   Ú	__class__s     €r!   rÍ   ÚCUDABackend.__init__¨   s   ø€ Ü‰Ñ˜Ô Ø!ˆ�r$   c                 ó  • SU;   a  US   S:X  a  SUS'   S[         R                  R                  =(       d    SU R                  R                   30nUR                  [        R                  R                  5        Vs0 s H  o3U;   d  M
  X   c  M  X1U   _M     sn5        [        U R                  US   5      5      nUR                  SS5      S:”  a  US	:  a  [        S
U S35      eSU;  aG  [        [        R                  5      nUS:¼  a  UR                  S5        [!        [#        U5      5      US'   SU;  a  US	:¼  a  SUS'   SU;  a  [         R$                  R&                  US'   US	:X  a  SOSUS'   [        S0 UD6$ s  snf )NrŠ   ÚconsanTr†   r'   Úsmrn   r   rA   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.r{   éY   Úfp8e4nvr|   )rz   ru   i   @r   r„   re   )r
   ÚruntimeÚoverride_archr   r'   Úupdaterj   Ú__dataclass_fields__Úkeysr%   rÄ   r‘   rÀ   Úsetr{   Úaddr•   r¥   ÚlanguageÚdefault_fp_fusion)r—   ÚoptsÚargsrž   r`   r{   s         r!   Úparse_optionsÚ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Ò)P AÐY]ÑT]“ZÐaeÑah“Z�Q˜Q™’ZÑ)PÑ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   Á<	FÂ	FÂ	Fc                 óH   • UR                   UR                  UR                  4$ rM   )rm   rn   Úshared)r—   Úmetadatas     r!   Úpack_metadataÚCUDABackend.pack_metadataË   s%   € à×ÑØ×ÑØ�O‰Oð
ð 	
r$   c                 óÎ   • SS K Js  Js  Jn  [	        U R                  UR                  5      5      nUS:¼  a  UR                  OUR                  [        U R                  5      S.nU$ )Nr   r;   )Úconvert_custom_typesr&   )Útriton.language.extra.cudarÝ   Úextrar‡   r%   rÄ   r'   Úconvert_custom_float8_sm80Úconvert_custom_float8_sm70r&   r   )r—   rO   r‡   r`   Úcodegen_fnss        r!   Úget_codegen_implementationÚ&CUDABackend.get_codegen_implementationÒ   sV   € ß1Ó1Ü˜×)Ñ)¨'¯,©,Ó7Ó8ˆ
ð 0:¸RÓ/?ˆD×+Ò+ÀT×EdÑEdÜ˜Ÿ™Ó%ñ
ˆð
 Ðr$   c                 ó   • SSK Jn  SU0$ )Nr   )r�   ztriton.language.extra.libdevice)rê   r�   )r—   r�   s     r!   Úget_module_mapÚCUDABackend.get_module_mapÜ   s   € Ý8Ø1°9Ð=Ð=r$   c                 óš   • [         R                  " U5        [        R                  (       a   [        R                  R                  U5        g g rM   )r	   Úload_dialectsr·   Úinstrumentation)r—   Úctxs     r!   rõ   ÚCUDABackend.load_dialectsà   s2   € Ü×Ò˜SÔ!Ü×&×&Ü×'Ñ'×5Ñ5°cÕ:ð 'r$   c                 óÊ  • [         R                  " U R                  5      nUR                  5         [        R
                  R                  U5        [        R                  R                  U5        US-  S:  a  [        R                  R                  U5        [        R
                  R                  U5        [        R                  R                  U5        [        R                  R                  U5        [        R
                  R                  U5        [        R
                  R                  U5        [        R                  R                  U5        UR!                  U S5        U $ )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û   ÚCUDABackend.make_ttirå   sà   € ä�_Š_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�‰×!Ñ! "Ô%Ü�‰×.Ñ.¨rÔ2Ø˜Ñ˜aÓÜ�K‰K×@Ñ@ÀÔDÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ Ô#Ü�‰×)Ñ)¨"Ô-Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ü�‰×#Ñ# BÔ'Ø
�‰ˆs�KÔ Øˆ
r$   c                 ó  • UR                   bI  U R                  S[        R                  " U R                  5      R                  UR                   5      5        [        R                  " U R                  5      nUR                  5       nUS-  S:¬  n[        R                  R                  USU 3UR                  SUR                  5        [        R                  R                  U5        [        R                  R                  XF5        [         R                  R"                  R%                  U5        [        R                  R'                  U5        [        R                  R)                  U5        [        R                  R+                  U5        [        R                  R'                  U5        [        R                  R-                  XCS:¬  5        [         R                  R"                  R/                  U5        [        R                  R1                  U5        US-  S;   GaC  [        R                  R3                  U5        [        R4                  R7                  U5        [        R                  R9                  U5        [        R4                  R7                  U5        [        R                  R;                  U5        [         R                  R<                  R?                  XBR@                  U5        [        R                  RC                  XBR@                  5        [        R                  RE                  U5        [        R                  RG                  XBR@                  U5        GOUS-  S:¼  Gaè  [        R                  R3                  U5        [        R4                  R7                  U5        [        R                  R9                  U5        [        R                  RI                  U5        [        R                  RK                  US5        [         R                  R"                  RM                  U5        [        R                  RC                  XBR@                  5        [        R                  RE                  U5        [        R                  RO                  XBR@                  5        [        R                  RG                  XBR@                  U5        [        R                  RQ                  U5        [        R                  R;                  U5        [        R                  RK                  US	5        [         R                  R"                  RS                  U5        O[        R                  R9                  U5        [        R4                  R7                  U5        [        R                  R1                  U5        [        R                  RU                  U5        [        R                  R-                  XCS:¬  5        [        R                  RW                  U5        [         R                  R"                  RY                  U5        US-  S
:¼  a)  [         R                  R"                  R[                  U5        [        R                  R'                  U5        [         R                  R"                  R]                  U5        [        R                  R_                  U5        [        R                  Ra                  U5        [        R                  R1                  U5        [        R4                  Rc                  U5        [         R                  R"                  Re                  XC5        [         R                  R"                  Rg                  U5        [        R4                  Ri                  U5        [        R4                  Rk                  U5        [        R4                  R7                  U5        URm                  U S5        U Ro                  5       US'   U $ )Nzttg.maxnregr>   r   rÇ   r   r;   )r   rú   FTrú   Ú
make_ttgirÚtensordesc_meta)8rr   Úset_attrr   Úbuilderrý   Úget_int32_attrrü   rþ   r   r  Úadd_convert_to_ttgpuirrm   rn   Ú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_warpspecrp   Ú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  Ú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°À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×/Ñ/°·N±NÔCÜ�N‰N×-Ñ-¨bÔ1Ü�N‰N×'Ñ'¨¯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×/Ñ/°·N±NÔCÜ�N‰N×-Ñ-¨bÔ1Ü�N‰N×.Ñ.¨r·>±>ÔBÜ�N‰N×'Ñ'¨¯N©N¸LÔIÜ�N‰N×7Ñ7¸Ô;Ü�N‰N×;Ñ;¸BÔ?ä�N‰N×/Ñ/°°DÔ9Ü�M‰M×#Ñ#×:Ñ:¸2Õ>ä�K‰K×'Ñ'¨Ô+Ü�‰×'Ñ'¨Ô+Ü�‰×&Ñ& rÔ*Ü�‰×#Ñ# BÔ'Ü�‰×0Ñ0°À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                 óö  • Un[         R                  " UR                  5      nUR                  5         [        R
                  R                  U5        [        R
                  R                  U5        [        R
                  R                  U5        [        R                  R                  R                  U5        [        R
                  R                  U5        [        R                  R                  U5        [        R                  R!                  U5        [        R
                  R                  U5        [        R"                  R%                  U5        UR'                  US5        UR)                  5       US'   U$ )NÚgluon_to_ttgirr  )r   rü   rý   rþ   r   Úgluonr   Úadd_infer_coalesced_encodingsÚadd_resolve_auto_encodingsr	   r  r2  r  rÿ   r8  r  r   r  r#  r
  r9  )r—   Úsrcrå   rO   r`   r  r  s          r!   r>  Ú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                 ó4  • [        X0R                  R                  5      nUn[        R                  " UR
                  5      nUR                  5         [        R                  R                  U5        [        R                  R                  U5        [        R                  R                  U5        [        R                  R                  U5        [        R                  R                  R!                  XtU5        [        R                  R"                  R%                  U5        [        R                  R"                  R'                  U5        [(        R*                  R,                  S:X  a  [        R                  R/                  U5        [        R                  R1                  U5        [        R                  R"                  R3                  Xt5        [4        R6                  (       a*  [4        R6                  R9                  SXvR
                  5        [        R                  R                  R;                  XtU5        [        R<                  R?                  U5        [        R<                  RA                  U5        [        R                  R"                  RC                  U5        [        R                  R"                  RE                  U5        [        R<                  R?                  U5        [        R<                  RA                  U5        [        R<                  RG                  U5        [        R                  RI                  U5        [(        R*                  RJ                  (       d>  [(        R*                  RL                  (       d  [        RN                  RQ                  U5        [4        R6                  (       a*  [4        R6                  R9                  SXvR
                  5        URS                  US5        [(        R*                  RL                  (       aá  [(        R*                  RJ                  (       da  [        R                  " UR
                  5      nUR                  5         [        RN                  RQ                  U5        URS                  US5        [        R                  " UR
                  5      nUR                  5         [        RN                  RU                  U5        URS                  US5        [V        RX                  " 5         [V        R
                  " 5       n[(        R*                  RZ                  (       a  []        S5      e[V        R^                  " Xh5      n	[a        U5      n
[c        X0R                  R                  5      nSn[        Rd                  " 5         [V        Rf                  " XœX«5        URh                  (       a  [        Rj                  " U	5        URl                  (       aQ  [        Rn                  " U	5      (       a6  URl                   VVs/ s H  u  pÞUPM	     nnn[V        Rp                  " XŸ5        [V        Rr                  " U	[V        Rt                  5        URw                  S	5      nUb  UUS
'   URw                  S5      US'   URw                  S5      US'   URw                  S5      US'   URw                  S5      US'   URw                  S5      =(       d    SUS'   URw                  S5      =(       d    SUS'   [y        U	5      nA	AU$ s  snnf )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-warpsrm   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)=rP   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_asanrF   Ú	to_modulerg   rV   Úset_short_ptrÚattach_datalayoutrv   Úset_nvvm_reflect_ftzr…   Úhas_extern_depsÚlink_extern_libsÚoptimize_moduleÚOPTIMIZE_O3Úget_int_attrrC   )r—   rB  rå   rO   r`   rN   r  r  rý   Úllvm_modÚprocrU   Útripler©   r2   ÚpathsÚtotal_num_warpsÚrets                     r!   rG  ÚCUDABackend.make_llirS  s·  € Ü2°7¿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¸"ÔIä×&×&Ü×'Ñ'×-Ñ-Ð.AÀ2Ç{Á{ÔSÜ�‰×Ñ×+Ñ+¨B¸KÔHÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×Ñ×1Ñ1°"Ô5Ü�‰×Ñ×;Ñ;¸BÔ?Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ü�‰×'Ñ'¨Ô+ä× Ñ ×2×2¼5×;LÑ;L×;o×;oÜ�M‰M×&Ñ& rÔ*ä×&×&Ü×'Ñ'×-Ñ-Ð.>ÀÇ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ä—>’> #Ó/ˆÜ& zÓ2ˆÜ ¯©×)9Ñ)9Ó:ˆØ&ˆÜ×ÒÔÜ×Ò˜x°Ô@Ø×%×%Ü×'Ò'¨Ô1à××¤6×#9Ò#9¸(×#CÑ#CØ.5×.AÒ.AÔBÒ.A™l˜t“TÑ.AˆEÑBÜ×!Ò! (Ô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           	      óÚ  • [        X0R                  R                  5      nSn[        U5      n[	        X0R                  R                  5      nS/n	[
        R                  " XXxX“R                  S5      n
[        R                  " SU
5      n[        U5      S:X  d   eUS   US'   US-   S	US-   3n[        R                  " S
SU 3U
[        R                  S9n
[        R                  " SSU 3U
[        R                  S9n
[        R                  R                  (       d  [        R                  " SSU
5      n
[        R                   R"                  (       a  [%        S5        [%        U
5        U
$ )NrH  znvptx-mad-wide-optFz(.visible .entry ([a-zA-Z_][a-zA-Z0-9_]*)r   r   r©   r>   r8   z\.version \d+\.\d+z	.version )Úflagsz\.target sm_\d+z.target sm_z,\s*debug|debug,\s*rc   z // -----// NVPTX Dump //----- //)rP   r   r'   rg   rV   r   Útranslate_to_asmru   r¾   ÚfindallÚlenÚsubÚ	MULTILINEr
   rT  r^  r	   Ú
dump_nvptxÚprint)r—   rB  rå   r  r`   rN   ro  rn  rU   ru  rr  Únamess               r!   Úmake_ptxÚCUDABackend.make_ptx³  s3  € Ü2°3¿¹×8HÑ8HÓIˆà&ˆÜ& zÓ2ˆÜ §[¡[×%5Ñ%5Ó6ˆØ%Ð&ˆÜ×#Ò# C°À×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           
      ó¼  • [        U R                  R                  5      R                  n[        R
                  " SSSS9 n[        R
                  " SSSS9 nUR                  U5        UR                  5         UR                  S-   n/ n	[        R                  R                  (       a  U	SS	/-  n	O,[        R                  R                  (       a  U	S
/-  n	OU	S/-  n	UR                  (       a  / OS/n
[        U5      n[        R                  R                  (       a  SS/O/ nUR                   (       a  UR                   R#                  S5      O/ nU/U	QU
QSPUQUQSU 3PUR                  PSPUPn [$        R&                  " USSUS9  [        R                  R(                  (       a7  [+        UR                  5       n[-        UR/                  5       5        S S S 5        [0        R                  R3                  UR                  5      (       a   [0        R4                  " UR                  5        [0        R                  R3                  UR                  5      (       a   [0        R4                  " UR                  5        [+        US5       nUR/                  5       nS S S 5        [0        R                  R3                  U5      (       a  [0        R4                  " U5        S S S 5        S S S 5        W$ ! , (       d  f       GN= f! [$        R6                   Ga  n[+        UR                  5       nUR/                  5       nS S S 5        O! , (       d  f       O= f[0        R                  R3                  UR                  5      (       a   [0        R4                  " UR                  5        UR8                  S:X  a  SnO3UR8                  S[:        R<                  -   :X  a  SnOSUR8                   3nU SW SSR?                  U5       S3n[-        SU SU S35        [A        U5      eS nAff = f! , (       d  f       GN•= f! , (       d  f       GNe= f! , (       d  f       W$ = f) NFÚwz.ptx)ÚdeleteÚmoderf   Ú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.
rX   )!r,   r   r'   r2   ÚtempfileÚNamedTemporaryFileÚwriteÚflushr©   r
   rT  r]  r	   Údisable_ptxas_optru   rg   rs   rE   r0   r
  Údump_ptxas_logrY   r|  r\   ÚosÚexistsÚremoveÚCalledProcessErrorÚ
returncodeÚsignalÚSIGSEGVr¦   r   )r—   rB  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_cubinÚCUDABackend.make_cubinÍ  s�  € Ü˜$Ÿ+™+×*Ñ*Ó+×0Ñ0ˆÜ×(Ò(°¸CÈÒOÐSWÜ×'Ò'¨u¸3ÀvÒNÐ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×#Ñ#Ü—	’	˜$”÷M O÷ Pð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 "Ö!ú÷G OÖNú÷ PÔOðP ˆús•   Á QÁDP:Å4AL Æ?K.ÇB$L É=P:Ê	P(ÊAP:ËQË.
K=Ë8L Ì P%ÌP Ì*M	Ì;	P Í
MÍCP Ð P%Ð%P:Ð(
P7Ð2P:Ð:
Q		ÑQÑ
Qc                 ó�  ^ ^^• T R                  TR                  5      mU[        R                  :X  a  UUU 4S jUS'   UUU 4S jUS'   OU[        R                  :X  a  UUU 4S jUS'   UUU 4S jUS'   UU 4S jUS	'   UU 4S
 jUS'   [
        R                  R                  b$  [
        R                  R                  T UTUT5        g g )Nc                 ó*   >• TR                  XTT5      $ rM   )rû   ©rB  rå   r`   rO   r—   s     €€€r!   Ú<lambda>Ú(CUDABackend.add_stages.<locals>.<lambda>  s   ø€ °4·>±>À#ÐQXÐZdÔ3er$   r  c                 ó*   >• TR                  XTT5      $ rM   )r  rª  s     €€€r!   r«  r¬    s   ø€ °D·O±OÀCÐSZÐ\fÔ4gr$   Úttgirc                 ó*   >• TR                  XTT5      $ rM   )r>  rª  s     €€€r!   r«  r¬    s   ø€ °D×4GÑ4GÈÐW^Ð`jÔ4kr$   c                 ó*   >• TR                  XTT5      $ rM   )rG  rª  s     €€€r!   r«  r¬     s   ø€ ¨t¯~©~¸cÈWÐV`Ô/ar$   Úllirc                 óR   >• TR                  XTTR                  R                  5      $ rM   )r~  r   r'   ©rB  rå   rO   r—   s     €€r!   r«  r¬  !  s   ø€ ¨d¯m©m¸CÈ7ÐTX×T_ÑT_×TdÑTdÔ.er$   Úptxc                 óR   >• TR                  XTTR                  R                  5      $ rM   )r¦  r   r'   r³  s     €€r!   r«  r¬  "  s   ø€ °·±ÀÈwÐX\×XcÑXc×XhÑXhÔ0ir$   rË   )rÄ   r'   r   ÚTRITONÚGLUONr
   rÖ   Úadd_stages_inspection_hook)r—   ÚstagesrO   rÝ   r`   s   ` ` @r!   Ú
add_stagesÚ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   • [        U R                  R                  5      nU SU R                  R                   3$ )Nr£   )r6   r   r'   )r—   r5   s     r!   r¬   ÚCUDABackend.hash&  s2   € ä# D§K¡K×$4Ñ$4Ó5ˆØ�˜!˜DŸK™K×,Ñ,Ð-Ð.Ð.r$   )rÎ   ) r®   r¯   r°   r±   rö   Ústaticmethodr   r»   rÄ   rC   rÈ   rÍ   r   rá   ræ   rï   r   r   rò   rõ   rû   r  r>  rG  r~  r¦  rº  Ú	functoolsÚ	lru_cacher¬   rµ   Ú__classcell__)rÏ   s   @r!   r·   r·   –   sÛ   ø† Ø€Oàð( 	ó (ó ð(ò#ð$¨#ô $ð"˜yð "¨T÷ "ð# Sô #ò>
òð>  S¨* _Ñ 5ô >ò;ð
 ñó ðð  ñGó ðGòRò&^ò@ò4JòXbð ×ÒÓñ/ó ö/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   rZ   r¾   r�  r˜  r“  r0   Úpathlibr   r&   r%   Ú
NvidiaToolr,   rÀ  r6   rK   rP   rV   r_   rg   rj   r·   re   r$   r!   Ú<module>rË     s6  ðß EÑ Eß 8Ó 8Ý Ý ,å !Û ß -Ó -Ý Û Û 	Û Û Û 	Û Ý ð#˜ô #ðO�Cð O˜E×,Ñ,ô Oð ×ÒÓñ˜Cô ó ðð ×ÒÓði Só ió ðið.°ô ð ×ÒÓð
 ó 
ó ð
ð ×Ò�TÓñ4ó ð4ð
&¨ô &ñ �$Ñ÷)?ð )?ó ð)?ôXS/�+õ S/r$   