Ë
    ÎÍ:jþM  ã                   óö   — 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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mZ d	efd
„Zd„ Zd„ Z ed¬«       G d„ d«      «       Z G d„ de«      Zy)é    )ÚBaseBackendÚ	GPUTargetÚLanguage)ÚirÚpassesÚllvmÚamd)Úknobs)Ú	dataclass)ÚAnyÚDictÚTuple)Ú
ModuleTypeN)ÚPathÚtargetc                 ó   — d„ S )Nc                  ó   — y)N©é   r   r   © )Úlhs_typeÚrhs_types     úq/home/mcse/projects/srt_converter/srt-converter-venv/lib/python3.12/site-packages/triton/backends/amd/compiler.pyú<lambda>z"get_min_dot_size.<locals>.<lambda>   s   � ó    r   ©r   s    r   Úget_min_dot_sizer      s
   € ñ 0Ð/r   c                 ót   — t         j                  j                  €| dk(  S t         j                  j                  S ©NÚgfx942)r
   r	   Úuse_block_pingpong©Úarchs    r   Úis_pingpong_schedule_enabledr$      s+   € Ü!&§¡×!=Ñ!=Ð!EˆD�HÑÐgÌ5Ï9É9×KgÑKgÐgr   c                 ót   — t         j                  j                  €| dk(  S t         j                  j                  S r   )r
   r	   Úuse_in_thread_transposer"   s    r   Úis_in_thread_transpose_enabledr'      s.   € Ü!&§¡×!BÑ!BÐ!JˆD�HÑÐqÔPU×PYÑPY×PqÑPqÐqr   T)Úfrozenc                   óR  — 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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eed#<   d$„ Zd%„ Z y)&Ú
HIPOptionsé   Ú	num_warpsr   Úwaves_per_eué   Ú
num_stagesÚnum_ctasNÚextern_libsr   Úcluster_dimsFÚdebugTÚsanitize_overflowr#   )Úfp8e5Úsupported_fp8_dtypesr   Ú!deprecated_fp8_dot_operand_dtypesÚieeeÚdefault_dot_input_precision)r8   Úallowed_dot_input_precisionsÚenable_fp_fusionÚlaunch_cooperative_gridr   Úmatrix_instr_nonkdimÚkpackÚallow_flush_denormÚmax_num_imprecise_acc_defaultÚhipÚbackend_nameÚnoneÚschedule_hintc                 ó8  — t        | j                  dd «      }|dk\  rdnd}t        j                  | d|«       | j                  dkD  r| j                  | j                  dz
  z  dk(  sJ d	«       ‚| j                  d
k(  r| j
                  dk(  sJ d«       ‚t        t        «      j                  dz  }| j                  €i nt        | j                  «      }dD ]  }t        ||› d�z  «      ||<   Œ t        j                  | dt        |j                  «       «      «       y )Né   éþÿÿÿé
   é    é@   Ú	warp_sizer   r   znum_warps must be a power of 2Úgfx950zgfx950 only accepts kpack == 1Úlib)ÚocmlÚocklz.bcr1   )Úintr#   ÚobjectÚ__setattr__r,   r>   r   Ú__file__Úparentr1   ÚdictÚstrÚtupleÚitems)ÚselfÚ	gfx_majorrK   Údefault_libdirr1   rM   s         r   Ú__post_init__zHIPOptions.__post_init__E   s  € Ü˜Ÿ	™	 ! B˜Ó(ˆ	Ø# rš/‘B¨rˆ	Ü×Ñ˜4 ¨iÔ8Ø�~‰~ Ò! t§~¡~¸¿¹È!Ñ9KÑ'LÐQRÒ&Rð 	0Ø/ó	0ÐRð �9‰9˜Ò Ø—:‘: ’?ÐDÐ$DÓD�?äœh›×.Ñ.°Ñ6ˆØ ×,Ñ,Ð4‘b¼$¸t×?OÑ?OÓ:PˆØ#ò 	AˆCÜ" >°s°e¸3°KÑ#?Ó@ˆK˜Òð	Aä×Ñ˜4 ´°k×6GÑ6GÓ6IÓ0JÕKr   c           	      óô   — dj                  | j                  j                  «       D ��cg c]  \  }}|› d|› �‘Œ c}}«      }t        j                  |j                  d«      «      j                  «       S c c}}w )NÚ_ú-úutf-8)ÚjoinÚ__dict__rX   ÚhashlibÚsha256ÚencodeÚ	hexdigest)rY   ÚnameÚvalÚkeys       r   ÚhashzHIPOptions.hashU   s]   € Ø�h‰h¸¿¹×9LÑ9LÓ9N×O©I¨D°#˜4˜&  # šÓOÓPˆÜ�~‰~˜cŸj™j¨Ó1Ó2×<Ñ<Ó>Ð>ùó Ps   ©A4
)!Ú__name__Ú
__module__Ú__qualname__r,   rP   Ú__annotations__r-   r/   r0   r1   rU   r2   rW   r3   Úboolr4   r#   rV   r6   r   r7   r9   r:   r;   r<   r=   r>   r?   r@   rB   rD   r\   rj   r   r   r   r*   r*      sö   … à€IˆsÓØ€L�#ÓØ€J�ÓØ€HˆcÓØ€K�ÓØ#€L�%Ó#Ø€Eˆ4ÓØ"Ð�tÓ"Ø€Dˆ#ÓØ'2Ð˜% ™*Ó2Ø46Ð% u¨S¡zÓ6Ø'-Ð Ó-Ø/9Ð  %¨¡*Ó9Ø!Ð�dÓ!Ø$)Ð˜TÓ)Ø !Ð˜#Ó!Ø€Eˆ3ƒNØ$Ð˜Ó$Ø)*Ð! 3Ó*Ø€L�#Óð"  €M�3ÓòLó ?r   r*   c                   óR  ‡ — e Zd Zedefd„«       Zdeddfˆ fd„Zde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ed„ «       Zed„ «       Zed„ «       Zed„ «       Zed„ «       Zed„ «       Zed„ «       Zed„ «       Zd„ Z ej:                  «       d„ «       Zˆ xZS )Ú
HIPBackendr   c                 ó    — | j                   dk(  S )NrA   )Úbackendr   s    r   Úsupports_targetzHIPBackend.supports_target\   s   € à�~‰~ Ñ&Ð&r   ÚreturnNc                 ój   •— t         ‰| �  |«       t        |j                  t        «      sJ ‚d| _        y )NÚhsaco)ÚsuperÚ__init__Ú
isinstancer#   rV   Ú
binary_ext)rY   r   Ú	__class__s     €r   ry   zHIPBackend.__init__`   s+   ø€ Ü‰Ñ˜Ô Ü˜&Ÿ+™+¤sÔ+Ð+Ð+Ø!ˆ�r   c                 ó    — d|j                   › �S )Núhip:r"   ©rY   Úoptionss     r   Úget_target_namezHIPBackend.get_target_namee   s   € Ø�g—l‘l�^Ð$Ð$r   c                 ó„  — dt         j                  j                  xs | j                  j                  i}| j                  j                  dk(  rBt        t        j                  «      }|j                  dh«       t        t        |«      «      |d<   d|vrµt        t        j                  «      }| j                  j                  dk(  r|j                  h d£«       nX| j                  j                  dk(  r|j                  dd	h«       n+d
| j                  j                  v r|j                  dd	h«       t        t        |«      «      |d<   d|vrt         j                  j                  |d<   |j                  t        j                  j                  «       D �ci c]  }||v r||   �|||   “Œ c}«       t        di |¤ŽS c c}w )Nr#   r    Útf32r:   r6   >   Úfp8e4b8Úfp8e4nvÚfp8e5b16rL   r…   r5   Úgfx12r;   r   )r
   ÚruntimeÚoverride_archr   r#   Úsetr*   r:   ÚupdaterW   Úsortedr6   ÚlanguageÚdefault_fp_fusionÚ__dataclass_fields__Úkeys)rY   ÚoptsÚargsr:   r6   Úks         r   Úparse_optionszHIPBackend.parse_optionsh   s�  € ØœŸ™×3Ñ3ÒG°t·{±{×7GÑ7GÐHˆð �;‰;×Ñ˜xÒ'Ü+.¬z×/VÑ/VÓ+WÐ(Ø(×/Ñ/°°Ô9Ü38¼Ð@\Ó9]Ó3^ˆDÐ/Ñ0à!¨Ñ-Ü#&¤z×'FÑ'FÓ#GÐ Ø�{‰{×Ñ 8Ò+Ø$×+Ñ+Ò,NÕOØ—‘×!Ñ! XÒ-Ø$×+Ñ+¨Y¸Ð,@ÕAØ˜DŸK™K×,Ñ,Ñ,Ø$×+Ñ+¨Y¸Ð,@ÔAÜ+0´Ð8LÓ1MÓ+NˆDÐ'Ñ(à TÑ)Ü',§~¡~×'GÑ'GˆDÐ#Ñ$Ø�‰¬×)HÑ)H×)MÑ)MÓ)Oö ; AØ˜T™	 d¨1¡gÐ&9ð ˜˜Q™‘Zò ;ô 	<äÑ!˜DÑ!Ð!ùò;s   ÆF=c                 óœ   — |j                   |j                  |j                  |j                  d   |j                  d   |j                  d   fS )Nr   r   r.   )r,   r0   Úsharedr2   )rY   Úmetadatas     r   Úpack_metadatazHIPBackend.pack_metadata�   sO   € à×ÑØ×ÑØ�O‰OØ×!Ñ! !Ñ$Ø×!Ñ! !Ñ$Ø×!Ñ! !Ñ$ð
ð 	
r   c                 ó0   — dt        | j                  «      iS )NÚmin_dot_size)r   r   r   s     r   Úget_codegen_implementationz%HIPBackend.get_codegen_implementation‹   s   € ØÔ 0°·±Ó =Ð>Ð>r   c                 ó   — ddl m} d|iS )Nr   )Ú	libdeviceztriton.language.extra.libdevice)Útriton.language.extra.hipr�   )rY   r�   s     r   Úget_module_mapzHIPBackend.get_module_mapŽ   s   € Ý7à1°9Ð=Ð=r   c                 ó.   — t        j                  |«       y ©N)r	   Úload_dialects)rY   Úctxs     r   r¢   zHIPBackend.load_dialects“   s   € Ü×Ñ˜#Õr   c                 óÔ   — dd l }d}t        | d«      r| j                  «       |k  S t        | |j                  «      r-t        | d«      r!| j                  «       j                  «       |k  S y)Nr   iÿÿÿÚ	ptr_rangeÚuntyped_storageF)ÚtorchÚhasattrr¥   rz   ÚTensorr¦   Úsize)Úargr§   Ú
MAX_INT_32s      r   Úis_within_2gbzHIPBackend.is_within_2gb–   s]   € ãàˆ
Ü�3˜Ô$Ø—=‘=“? jÑ0Ð0Ü�c˜5Ÿ<™<Ô(¬W°SÐ:KÔ-LØ×&Ñ&Ó(×-Ñ-Ó/°:Ñ=Ð=Ør   c                 óH   — t        j                  | «      }d| v r|ddggz  }|S )NÚSztt.pointer_rangerI   )r   Ú
parse_attr)ÚdescÚrets     r   r°   zHIPBackend.parse_attr¡   s1   € ä×$Ñ$ TÓ*ˆØ�$‰;ØÐ'¨Ð,Ð-Ñ-ˆCØˆ
r   c                 ó¦   — t        j                  | |fi |¤Ž}t        j                  j                  r|dk(  rt
        j                  | «      r|dz  }|S )NÚtensorr¯   )r   Úget_arg_specializationr
   r	   Úuse_buffer_opsrq   r­   )r«   ÚtyÚkwargsr²   s       r   rµ   z!HIPBackend.get_arg_specialization¨   sJ   € ä×0Ñ0°°bÑC¸FÑCˆô �9‰9×#Ò#¨¨hª¼:×;SÑ;SÐTWÔ;XØ�3‰JˆCØˆ
r   c                  óZ  — t         j                  j                  } | �t        | «      }|j	                  «       r|S t        t
        «      j                  dz  }|j	                  «       r|S t        d«      }|j	                  «       r|S t        d«      }|j	                  «       r|S t        d«      ‚)Nzllvm/bin/ld.lldz/opt/rocm/llvm/bin/ld.lldz/usr/bin/ld.lldzWROCm linker /opt/rocm/llvm/bin/ld.lld not found. Set 'TRITON_HIP_LLD_PATH' to its path.)r
   r	   Úlld_pathr   Úis_filerS   rT   Ú	Exception)Úlld_env_pathÚllds     r   Úpath_to_rocm_lldzHIPBackend.path_to_rocm_lld±   s’   € ô —y‘y×)Ñ)ˆØÐ#Ü�|Ó$ˆCØ�{‰{Œ}Ø�
ä”8‹n×#Ñ#Ð&7Ñ7ˆØ�;‰;Œ=ØˆJÜÐ.Ó/ˆØ�;‰;Œ=ØˆJÜÐ$Ó%ˆØ�;‰;Œ=ØˆJÜÐqÓrÐrr   c                 óò  — t        j                  | j                  «      }|j                  «        t        j
                  j                  |«       t        j                  j                  |«       t        j                  j                  |«       t        j
                  j                  |«       t        j                  j                  |«       t        j                  j                  |«       t        j
                  j                  |«       t        j                  j                  |«       t        j
                  j                  |«       t        j                  j!                  |«       |j#                  | «       | S r¡   )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_triton_licmÚadd_symbol_dceÚadd_loop_unrollÚrun)Úmodr—   r€   Úpms       r   Ú	make_ttirzHIPBackend.make_ttirÅ   sâ   € ä�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�‰×!Ñ! "Ô%Ü�‰×.Ñ.¨rÔ2Ü�‰×<Ñ<¸RÔ@Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ Ô#Ü�‰×)Ñ)¨"Ô-Ü�‰×Ñ˜bÔ!Ü�‰×#Ñ# BÔ'Ü�‰×$Ñ$ RÔ(Ü�‰×#Ñ# BÔ'Ø
�‰ˆsŒØˆ
r   c                 ó¦  — t        j                  | j                  «      }|j                  «        t        j
                  j                  |d|j                  › �|j                  |j                  |j                  «       |j                  | «       t        j                  | j                  «      }|j                  «        t        j                  j                  |«       t        j                  j                  |«       t        j                  j                  |«       t         j                  j                  j#                  ||j                  |j$                  |j&                  «       t        j                  j                  |«       t         j                  j                  j)                  |«       t        j                  j+                  |d«       t         j                  j                  j-                  |«       t        j                  j/                  |«       t        j0                  j3                  |«       t        j
                  j5                  |«       t        j0                  j3                  |«       t6        j                   j8                  }t6        j                   j:                  }t6        j                   j<                  }|j>                  dk(  rdx}}t         j                  j                  jA                  ||jB                  |||«       |r4t         j                  j                  jE                  ||j                  «       t        j0                  j3                  |«       |j>                  jG                  «       dk7  r4t         j                  j                  jI                  ||j>                  «       t        j                  j+                  |d«       t        j                  j                  |«       t        j                  jK                  |«       tM        |j                  «      rHt         j                  j                  jO                  |«       t        j                  j                  |«       t         j                  j                  jQ                  |«       tS        |j                  «      }|rC|jB                  dk(  r4t         j                  j                  jU                  ||jB                  «       t6        j                   jV                  r|t         j                  j                  jY                  |«       t        j0                  j3                  |«       t         j                  j                  j[                  ||j                  «       t         j                  j                  j]                  |«       t        j0                  j3                  |«       t        j0                  j_                  |«       t        j0                  ja                  |«       |r4t         j                  j                  jc                  ||j                  «       |j                  | «       | S )Nr~   Tzlocal-prefetchr   rC   r.   )2r   rÁ   rÂ   rÃ   r   rÆ   Úadd_convert_to_ttgpuirr#   r,   rK   r0   rÐ   ÚttgpuirÚadd_coalesceÚadd_remove_layout_conversionsÚadd_optimize_thread_localityr	   Úadd_accelerate_matmulr=   r>   Úadd_optimize_epilogueÚadd_optimize_dot_operandsÚadd_hoist_layout_conversionsÚadd_fuse_nested_loopsrÄ   rÉ   rÍ   r
   Úglobal_prefetchÚlocal_prefetchÚuse_async_copyrD   Úadd_stream_pipeliner/   Úadd_coalesce_async_copyÚlowerÚinsert_instruction_sched_hintsÚadd_reduce_data_duplicationr'   Úadd_in_thread_transposeÚadd_reorder_instructionsr$   Úadd_block_pingpongr¶   Úadd_canonicalize_pointersÚadd_convert_to_buffer_opsÚadd_fold_true_cmpirÌ   rÎ   Úadd_update_async_wait_count)rÑ   r—   r€   rÒ   rß   rà   rá   r!   s           r   Ú
make_ttgirzHIPBackend.make_ttgirÖ   sö  € ä�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�‰×*Ñ*¨2°°g·l±l°^Ð/DÀg×FWÑFWÐY`×YjÑYjØ+2×+;Ñ+;ô	=à
�‰ˆsŒÜ�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�‰×#Ñ# BÔ'Ü�‰×4Ñ4°RÔ8Ü�‰×3Ñ3°BÔ7Ü�
‰
×Ñ×0Ñ0°°W·\±\À7×C_ÑC_Ðah×anÑanÔoÜ�‰×4Ñ4°RÔ8Ü�
‰
×Ñ×0Ñ0°Ô4Ü�‰×0Ñ0°°TÔ:Ü�
‰
×Ñ×7Ñ7¸Ô;ä�‰×,Ñ,¨RÔ0Ü�‰×'Ñ'¨Ô+Ü�‰×#Ñ# BÔ'Ü�‰×'Ñ'¨Ô+äŸ)™)×3Ñ3ˆÜŸ™×1Ñ1ˆÜŸ™×1Ñ1ˆð × Ñ Ð$4Ò4Ø/0Ð0ˆO˜nä�
‰
×Ñ×.Ñ.¨r°7×3EÑ3EÀÐXfÐhvÔwÙÜ�J‰J×Ñ×6Ñ6°r¸7¿<¹<ÔHÜ�‰×'Ñ'¨Ô+Ø× Ñ ×&Ñ&Ó(¨FÒ2Ü�J‰J×Ñ×=Ñ=¸bÀ'×BWÑBWÔXÜ�‰×0Ñ0°°TÔ:Ü�‰×4Ñ4°RÔ8Ü�‰×2Ñ2°2Ô6Ü)¨'¯,©,Ô7Ü�J‰J×Ñ×6Ñ6°rÔ:Ü�N‰N×8Ñ8¸Ô<Ü�
‰
×Ñ×3Ñ3°BÔ7Ü9¸'¿,¹,ÓGÐÙ '×"4Ñ"4¸Ò"9Ü�J‰J×Ñ×1Ñ1°"°g×6HÑ6HÔIä�9‰9×#Ò#Ü�J‰J×Ñ×8Ñ8¸Ô<Ü�M‰M×+Ñ+¨BÔ/Ü�J‰J×Ñ×8Ñ8¸¸W¿\¹\ÔJä�
‰
×Ñ×-Ñ-¨bÔ1Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(ÙÜ�J‰J×Ñ×:Ñ:¸2¸w¿|¹|ÔLØ
�‰ˆsŒØˆ
r   c                 óÀ  — | }t        j                  |j                  «      }|j                  «        t        j
                  j                  |«       t        j                  j                  |«       t        j                  j                  |«       t        j
                  j                  |«       t        j
                  j                  |«       |j                  |«       |S r¡   )r   rÁ   rÂ   rÃ   r   rÖ   rÅ   rÄ   Úadd_sccprÆ   Úadd_loop_aware_cserÉ   Ú add_combine_tensor_select_and_ifrÐ   )Úsrcr—   r€   rÑ   rÒ   s        r   Ú	ttgir_optzHIPBackend.ttgir_opt  s�   € àˆÜ�_‰_˜SŸ[™[Ó)ˆØ
�‰Ôä�‰×"Ñ" 2Ô&Ü�‰×Ñ˜rÔ"Ü�‰×&Ñ& rÔ*Ü�‰×(Ñ(¨Ô,Ü�‰×7Ñ7¸Ô;à
�‰ˆsŒØˆ
r   c                 ó$  — | }t        j                  |j                  «      }|j                  «        d}t        j
                  j                  j                  ||j                  |«       t
        j                  j                  |«       t
        j                  j                  |«       t
        j                  j                  |«       d}t        j
                  j                  j                  ||j                  |«       t
        j                  j                  |«       t
        j                  j!                  |«       t
        j                  j#                  |«       t
        j                  j%                  |«       t
        j                  j                  |«       t
        j                  j!                  |«       t
        j                  j'                  |«       |j(                  j+                  «       dk7  r?t        j
                  j                  j-                  ||j                  |j.                  «       t0        j2                  j4                  st
        j6                  j9                  |«       t        j
                  j                  j;                  ||«       |j=                  |«       t?        j@                  «        t?        j                  «       }t?        jB                  ||«      }t	        jD                  |«       d}	t0        j2                  jF                  rd}	t?        jH                  |t        jJ                  |j                  |	«       t	        jL                  ||j                  «       t	        jN                  |d«       t	        jP                  |dd«       t	        jP                  |d	d«       t	        jP                  |d
d«       t	        jP                  |d|jR                  dk(  «       |jU                  «       D �
cg c]  }
|
jW                  «       rŒ|
‘Œ }}
|d   jY                  t        jZ                  «       |d   j]                  dd|j^                  |jR                  z  › �«       |d   j]                  d|j`                  › «       |jb                  rdnd}|d   j]                  d|«       t0        j2                  jF                  r'|d   je                  d«       |d   jg                  «        t	        jh                  |d   «       t0        j2                  jF                  r\tk        tl        «      jn                  dz  }tq        |dz  «      tq        |dz  «      tq        |dz  «      g}t?        jr                  ||«       nW|jt                  rK|jt                  D ��cg c]  \  }}t	        jv                  ||«      sŒ|‘Œ  }}}t?        jr                  ||«       t?        jx                  |t>        jz                  |j                  dg |j|                  «       t0        j                  j~                  rt	        j€                  |d   «       | jƒ                  d«      |d<   t	        j„                  |«       t	        j†                  |«       tq        |«      S c c}
w c c}}w )Nr   TrC   Ú ú+xnackiô  Ú__oclc_finite_only_optFÚ__oclc_correctly_rounded_sqrt32Ú__oclc_unsafe_math_optÚ__oclc_wavefrontsize64rJ   zamdgpu-flat-work-group-sizez1,zamdgpu-waves-per-euzpreserve-signr8   zdenormal-fp-math-f32rM   z
asanrtl.bczocml.bczockl.bcz
ttg.sharedr–   )Dr   rÁ   rÂ   rÃ   r	   r   rÖ   Úadd_optimize_lds_usager#   ÚconvertÚadd_scf_to_cfÚadd_index_to_llvmirÚadd_allocate_shared_memoryÚadd_to_llvmirrÄ   rÉ   rÌ   Úadd_cf_to_llvmirÚadd_arith_to_llvmirrÎ   rD   rä   Úlower_instruction_sched_hintsr/   r
   ÚcompilationÚdisable_line_infoÚllvmirÚadd_di_scopeÚadd_builtin_func_to_llvmirrÐ   r   Úinit_targetsÚ	to_moduleÚattach_target_tripleÚenable_asanÚattach_datalayoutÚTARGET_TRIPLEÚset_isa_versionÚset_abi_versionÚset_bool_control_constantrK   Úget_functionsÚis_declarationÚset_calling_convÚCALLING_CONV_AMDGPU_KERNELÚadd_fn_attrr,   r-   r?   Úadd_fn_target_featureÚadd_fn_asan_attrÚset_all_fn_arg_inregr   rS   rT   rV   Úlink_extern_libsr1   Úneed_extern_libÚoptimize_moduleÚOPTIMIZE_O3r;   Úscalarize_packed_fopsÚ#add_scalarize_packed_fops_llvm_passÚget_int_attrÚcleanup_bitcode_metadataÚdisable_print_inline)ró   r—   r€   rÑ   rÒ   Úcustom_lds_sizeÚ_HIPBackend__HIP_FTZrÂ   Úllvm_modÚtarget_featuresÚfnÚfnsÚdenormal_moder[   Úpathsrg   Úpaths                    r   Ú	make_llirzHIPBackend.make_llir#  sÌ  € àˆä�_‰_˜SŸ[™[Ó)ˆØ
�‰Ôð ˆÜ�
‰
×Ñ×1Ñ1°"°g·l±lÀOÔTÜ�‰×$Ñ$ RÔ(Ü�‰×*Ñ*¨2Ô.ä�‰×1Ñ1°"Ô5ð ˆ	Ü�
‰
×Ñ×(Ñ(¨¨W¯\©\¸9ÔEÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!ä�‰×'Ñ'¨Ô+Ü�‰×*Ñ*¨2Ô.Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ø× Ñ ×&Ñ&Ó(¨FÒ2Ü�J‰J×Ñ×<Ñ<¸RÀÇÁÈw×OaÑOaÔbÜ× Ñ ×2Ò2Ü�M‰M×&Ñ& rÔ*Ü�
‰
×Ñ×5Ñ5°b¸)ÔDØ
�‰ˆsŒô 	×ÑÔÜ—,‘,“.ˆÜ—>‘> # wÓ/ˆÜ× Ñ  Ô*ØˆÜ×Ñ×(Ò(Ø&ˆOÜ×Ñ˜x¬×):Ñ):¸G¿L¹LÈ/ÔZô 	×Ñ˜H g§l¡lÔ3Ü×Ñ˜H cÔ*Ü×%Ñ% hÐ0HÈ%ÔPÜ×%Ñ% hÐ0QÐSWÔXÜ×%Ñ% hÐ0HÈ%ÔPÜ×%Ñ% hÐ0HÈ'×J[ÑJ[Ð_aÑJaÔbð %×2Ñ2Ó4ÖP�b¸B×<MÑ<MÕ<OŠrÐPˆÐPàˆA‰×Ñ¤× >Ñ >Ô?ØˆA‰×ÑÐ8¸B¸w×?PÑ?PÐQX×QbÑQbÑ?bÐ>cÐ:dÔeð 	ˆA‰×ÑÐ0°W×5IÑ5IÐ4JÔLØ+2×+EÒ+E™È6ˆØˆA‰×ÑÐ1°=ÔAÜ×Ñ×(Ò(Ø�‰F×(Ñ(¨Ô2Ø�‰F×#Ñ#Ô%ô
 	× Ñ   Q¡Ô(ä×Ñ×(Ò(Ü!¤(›^×2Ñ2°UÑ:ˆNä�N \Ñ1Ó2Ü�N YÑ.Ó/Ü�N YÑ.Ó/ðˆEô
 ×!Ñ! (¨EÕ2Ø× Ò Ø.5×.AÑ.A×i™l˜t TÄS×EXÑEXÐYaÐcgÕEh’TÐiˆEÑiÜ×!Ñ! (¨EÔ2ä×Ñ˜X¤t×'7Ñ'7¸¿¹ÀrÈ2Èw×OgÑOgÔhä�9‰9×*Ò*Ü×3Ñ3°C¸±FÔ;ð !×-Ñ-¨lÓ;ˆ�Ñä×$Ñ$ XÔ.ô 	× Ñ  Ô*Ü�8‹}Ðùò_ Qùó@ js   Ï8ZÐZÖZÖ3Zc           	      ó€  — t        j                  d| «      }t        |«      dk(  sJ ‚|d   |d<   g }|j                  dk(  r|j	                  d«       t        j                  | t        j                  |j                  d||j                  d«      }t        j                  j                  rt        d	«       t        |«       |S )
Nz3define amdgpu_kernel void @([a-zA-Z_][a-zA-Z0-9_]*)r   r   rg   Ú	attentionzsink-insts-to-avoid-spillsrö   Fz!// -----// AMDGCN Dump //----- //)ÚreÚfindallÚlenrD   Úappendr   Útranslate_to_asmr	   r  r#   r;   r
   Údump_amdgcnÚprint)ró   r—   r€   ÚnamesÚflagsÚamdgcns         r   Úmake_amdgcnzHIPBackend.make_amdgcn‘  s§   € ô
 —
‘
ÐQÐSVÓWˆÜ�5‹z˜QŠÐˆØ  ™8ˆ�Ñàˆð
 × Ñ  KÒ/Ø�L‰LÐ5Ô6Ü×&Ñ& s¬C×,=Ñ,=¸w¿|¹|ÈRÐQVÐX_×XpÑXpÐrwÓxˆÜ�9‰9× Ò ÜÐ5Ô6Ü�&ŒMØˆr   c                 ó   — d}t         j                  j                  rd}t        j                  | |j
                  |«      }t        j                  «       }t        j                  «       5 }t        j                  «       5 }t        |j                  d«      5 }|j                  |«       d d d «       t        j                  |ddd|j                  d|j                  g«       d d d «       t        |j                  d«      5 }	|	j                  «       }
d d d «       d d d «       
S # 1 sw Y   ŒzxY w# 1 sw Y   ŒNxY w# 1 sw Y   Œ+xY w# 1 sw Y   
S xY w)	Nrö   r÷   Úwbz-flavorÚgnuz-sharedz-oÚrb)r
   r  r  r	   Úassemble_amdgcnr#   rq   r¿   ÚtempfileÚNamedTemporaryFileÚopenrg   ÚwriteÚ
subprocessÚ
check_callÚread)ró   r—   r€   r'  rw   Ú	rocm_pathÚtmp_outÚtmp_inÚfd_inÚfd_outr²   s              r   Ú
make_hsacozHIPBackend.make_hsaco§  s   € àˆÜ×Ñ×(Ò(Ø&ˆOÜ×#Ñ# C¨¯©°ÓGˆä×/Ñ/Ó1ˆ	Ü×(Ñ(Ó*ð 	$¨gÜ×,Ñ,Ó.ð q°&Ü˜&Ÿ+™+ tÓ,ð '°Ø—K‘K Ô&÷'ä×%Ñ% y°)¸UÀIÈvÏ{É{Ð\`Ðbi×bnÑbnÐ&oÔp÷qô �g—l‘l DÓ)ð $¨VØ—k‘k“m�÷$÷	$ð ˆ
÷'ð 'ú÷qð qú÷$ð $ú÷	$ð ˆ
úsT   Á(EÁ=D+ÂDÂ&8D+ÃEÃ<D7ÄEÄD(Ä$D+Ä+D4	Ä0EÄ7E 	Ä<EÅEc                 óÂ   ‡ ‡— |t         j                  k(  rˆˆ fd„|d<   ˆˆ fd„|d<   n|t         j                  k(  r	ˆˆ fd„|d<   ˆˆ fd„|d<   ˆˆ fd„|d	<   ˆˆ fd
„|d<   y )Nc                 ó*   •— ‰j                  | |‰«      S r¡   )rÓ   ©ró   r—   r€   rY   s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>º  s   ø€ °4·>±>À#ÀxÐQXÓ3Y€ r   rÆ   c                 ó*   •— ‰j                  | |‰«      S r¡   )rî   rO  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>»  s   ø€ °D·O±OÀCÈÐSZÓ4[€ r   Úttgirc                 ó*   •— ‰j                  | |‰«      S r¡   )rô   rO  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>½  s   ø€ °D·N±NÀ3ÈÐRYÓ4Z€ r   c                 ó*   •— ‰j                  | |‰«      S r¡   )r-  rO  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>¾  s   ø€ ¨t¯~©~¸cÀ8ÈWÓ/U€ r   Úllirc                 ó*   •— ‰j                  | |‰«      S r¡   )r:  rO  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>¿  s   ø€ °×1AÑ1AÀ#ÀxÐQXÓ1Y€ r   r9  c                 ó*   •— ‰j                  | |‰«      S r¡   )rL  rO  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>À  s   ø€ °·±ÀÀXÈwÓ0W€ r   rw   )r   ÚTRITONÚGLUON)rY   Ústagesr€   r�   s   ` ` r   Ú
add_stageszHIPBackend.add_stages¸  sR   ù€ Ø”x—‘Ò&ÜYˆF�6‰NÜ[ˆF�7ŠOØœŸ™Ò'ÜZˆF�7‰OÜUˆˆv‰ÜYˆˆxÑÜWˆˆwŠr   c                 óz   — t        j                  t        j                  «       dgd¬«      }|› d| j                  › �S )Nz	--versionr`   )Úencodingr_   )rD  Úcheck_outputrq   r¿   r   )rY   Úversions     r   rj   zHIPBackend.hashÂ  s8   € ä×)Ñ)¬:×+FÑ+FÓ+HÈ+Ð*VÐahÔiˆØ�˜!˜DŸK™K˜=Ð)Ð)r   ) rk   rl   rm   Ústaticmethodr   rt   ry   rV   r�   r   r”   r˜   r›   r   r   rŸ   r¢   r­   r°   rµ   r¿   rÓ   rî   rô   r-  r:  rL  rZ  Ú	functoolsÚ	lru_cacherj   Ú__classcell__)r|   s   @r   rq   rq   Z   s^  ø„ àð' 	ò 'ó ð'ð"˜yð "¨Tõ "ð
%¨#ó %ð" Só "ò2
ò?ð>  S¨* _Ñ 5ó >ò
ð ñó ðð ñó ðð ñó ðð ñsó ðsð& ñó ðð  ñ;ó ð;ðz ñó ðð ñkó ðkðZ ñó ðð* ñó ðò Xð €Y×ÑÓñ*ó ô*r   rq   )Útriton.backends.compilerr   r   r   Útriton._C.libtritonr   r   r   r	   Útritonr
   Údataclassesr   Útypingr   r   r   Útypesr   rc   r@  r0  rD  r`  Úpathlibr   r   r$   r'   r*   rq   r   r   r   ú<module>rj     sw   ðß EÑ Eß 5Ó 5Ý Ý !ß #Ñ #Ý Û Û Û 	Û Û Ý ð0˜Yó 0òhòrñ �$Ô÷9?ð 9?ó ð9?ôxk*�õ k*r   