Ë
    0²j~^  ã                   ó  — 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Zd dlmZ d	efd
„Z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     ú[/var/www/html/ia-enhanced/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         j                  j                  €| dk(  xs | dk(  xr |du S t         j                  j                  S )NÚgfx942Úgfx950T)r
   r	   Úuse_block_pingpong)ÚarchÚuse_async_copys     r   Úis_pingpong_schedule_enabledr#      sI   € ä�9‰9×'Ñ'Ð/ð �HÑÒM ¨Ñ!1Ò!L°nÈÐ6Lð RÜ5:·Y±Y×5QÑ5QðRr   c                 ót   — t         j                  j                  €| dk(  S t         j                  j                  S )Nr   )r
   r	   Úuse_in_thread_transpose©r!   s    r   Úis_in_thread_transpose_enabledr'      s.   € Ü!&§¡×!BÑ!BÐ!JˆD�HÑÐqÔPU×PYÑPY×PqÑPqÐqr   c                 ó6   — t         j                  j                  S ©N)r
   r	   r"   r&   s    r   Úis_async_copy_enabledr*      s   € Ü�9‰9×#Ñ#Ð#r   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   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eed#<   d$„ Zd%„ Zy	)&Ú
HIPOptionsé   Ú	num_warpsr   Úwaves_per_eué   Ú
num_stagesr   Únum_ctasNÚextern_libsFÚdebugTÚsanitize_overflowr!   )Úfp8e4nvÚfp8e5Úfp8e5b16Úfp8e4b8Úsupported_fp8_dtypesr   Ú!deprecated_fp8_dot_operand_dtypesÚieeeÚdefault_dot_input_precision)r=   Úbf16x3Úbf16x6Úallowed_dot_input_precisionsÚenable_fp_fusionÚlaunch_cooperative_gridÚmatrix_instr_nonkdimÚkpackÚallow_flush_denormÚmax_num_imprecise_acc_defaultÚhipÚbackend_nameÚ Úinstrumentation_modeÚnoneÚschedule_hintc                 óž  — 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(  rI| j
                  dk7  r:t        j                  d| j
                  › d�«       t        j                  | d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 2r   zckpack is deprecated starting from gfx950 and will be removed in later releases. So for now kpack = z7 will be overwritten to 1 to make transitioning easier.rE   Úlib)ÚocmlÚocklz.bcr4   )Úintr!   ÚobjectÚ__setattr__r/   rE   ÚwarningsÚwarnr   Ú__file__Úparentr4   ÚdictÚstrÚtupleÚitems)ÚselfÚ	gfx_majorrT   Údefault_libdirr4   rU   s         r   Ú__post_init__zHIPOptions.__post_init__S   sG  € Ü˜Ÿ	™	 ! B˜Ó(ˆ	Ø# rš/‘B¨rˆ	Ü×Ñ˜4 ¨iÔ8Ø�~‰~ Ò! t§~¡~¸¿¹È!Ñ9KÑ'LÐQRÒ&Rð 	-Ø,ó	-ÐRð �I‰I˜Ò!¨¯
©
°aªÜ�M‰MØuÐvz÷  wAñ  wAð  vBð  Byð  zôô ×Ñ˜t W¨aÔ0äœ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__rb   ÚhashlibÚsha256ÚencodeÚ	hexdigest)rc   ÚnameÚvalÚkeys       r   ÚhashzHIPOptions.hashf   s]   € Ø�h‰h¸¿¹×9LÑ9LÓ9N×O©I¨D°#˜4˜&  # šÓOÓPˆÜ�~‰~˜cŸj™j¨Ó1Ó2×<Ñ<Ó>Ð>ùó Ps   ©A4
) Ú__name__Ú
__module__Ú__qualname__r/   rX   Ú__annotations__r0   r2   r3   r4   r_   r5   Úboolr6   r!   r`   r;   r   r<   r>   rA   rB   rC   rD   rE   rF   rG   rI   rK   rM   rf   rt   r   r   r   r-   r-   #   sú   … à€IˆsÓØ€L�#ÓØ€J�ÓØ€HˆcÓØ€K�ÓØ€Eˆ4ÓØ"Ð�tÓ"Ø€Dˆ#Óð
 (SÐ˜% ™*ÓRØ46Ð% u¨S¡zÓ6Ø'-Ð Ó-Ø/KÐ  %¨¡*ÓKØ!Ð�dÓ!Ø$)Ð˜TÓ)Ø !Ð˜#Ó!Ø€Eˆ3ƒNØ$Ð˜Ó$Ø)*Ð! 3Ó*Ø€L�#ÓØ "Ð˜#Ó"ð*  €M�3ÓòLó&?r   r-   c                   óJ  ‡ — e Zd Zd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d„ Z ej<                  «       d„ «       Zˆ xZ S )Ú
HIPBackendNFr   c                 ó    — | j                   dk(  S )NrH   )Úbackendr   s    r   Úsupports_targetzHIPBackend.supports_targeto   s   € à�~‰~ Ñ&Ð&r   Úreturnc                 ój   •— t         ‰| �  |«       t        |j                  t        «      sJ ‚d| _        y )NÚhsaco)ÚsuperÚ__init__Ú
isinstancer!   r`   Ú
binary_ext)rc   r   Ú	__class__s     €r   rƒ   zHIPBackend.__init__s   s+   ø€ Ü‰Ñ˜Ô Ü˜&Ÿ+™+¤sÔ+Ð+Ð+Ø!ˆ�r   c                 ó    — d|j                   › �S )Núhip:r&   ©rc   Úoptionss     r   Úget_target_namezHIPBackend.get_target_namex   s   € Ø�g—l‘l�^Ð$Ð$r   c                 óà  — dt         j                  j                  xs | j                  j                  i}|j                  dd«      dkD  rKt        j                  | j                  j                  «      s"t        d| j                  j                  › �«      ‚| j                  j                  dk(  rBt        t        j                  «      }|j                  dh«       t        t        |«      «      |d<   d|vr%t        t        t        j                  «      «      |d<   | j                  j                  d	k(  rCt        t        j                   «      }|j                  d
dh«       t        t        |«      «      |d<   d|vrt         j"                  j$                  |d<   |j                  t        j&                  j)                  «       D �ci c]  }||v sŒ||   €Œ|||   “Œ c}«       t        di |¤ŽS c c}w )Nr!   r3   r   znum_ctas > 1 not supported on r   Útf32rA   r;   r   r9   r:   r<   rB   r   )r
   ÚruntimeÚoverride_archr   r!   Úgetr	   Úsupports_multi_cta_launchÚ
ValueErrorÚsetr-   rA   Úupdatera   Úsortedr;   r<   ÚlanguageÚdefault_fp_fusionÚ__dataclass_fields__Úkeys)rc   ÚoptsÚargsrA   r<   Úks         r   Úparse_optionszHIPBackend.parse_options{   s›  € ØœŸ™×3Ñ3ÒG°t·{±{×7GÑ7GÐHˆà�8‰8�J Ó" QÒ&¬s×/LÑ/LÈTÏ[É[×M]ÑM]Ô/^ÜÐ=¸d¿k¹k×>NÑ>NÐ=OÐPÓQÐQð �;‰;×Ñ˜xÒ'Ü+.¬z×/VÑ/VÓ+WÐ(Ø(×/Ñ/°°Ô9Ü38¼Ð@\Ó9]Ó3^ˆDÐ/Ñ0à!¨Ñ-Ü+0´¼
×8WÑ8WÓ1XÓ+YˆDÐ'Ñ(à�;‰;×Ñ˜xÒ'Ü03´J×4`Ñ4`Ó0aÐ-Ø-×4Ñ4°jÀ)Ð5LÔMÜ8=¼fÐEfÓ>gÓ8hˆDÐ4Ñ5à TÑ)Ü',§~¡~×'GÑ'GˆDÐ#Ñ$Ø�‰¬×)HÑ)H×)MÑ)MÓ)OÖu AÐSTÐX\ÒS\ÐaeÐfgÑahÑat�Q˜˜Q™‘ZÒuÔvÜÑ!˜DÑ!Ð!ùò vs   Ç	G+ÇG+ÇG+c                 óH   — |j                   |j                  |j                  fS r)   )r/   r3   Úshared)rc   Úmetadatas     r   Úpack_metadatazHIPBackend.pack_metadata”   s%   € à×ÑØ×ÑØ�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¦   )rc   r¦   s     r   Úget_module_mapzHIPBackend.get_module_mapž   s   € Ý7à1°9Ð=Ð=r   c                 óŽ   — t        j                  |«       t        j                  r t        j                  j                  |«       y y r)   )r	   Úload_dialectsr{   Úinstrumentation)rc   Úctxs     r   rª   zHIPBackend.load_dialects£   s2   € Ü×Ñ˜#ÔÜ×%Ò%Ü×&Ñ&×4Ñ4°SÕ9ð &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®   r„   Ú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_rangerR   )r   Ú
parse_attr)ÚdescÚrets     r   r¹   zHIPBackend.parse_attr³   s1   € ä×$Ñ$ TÓ*ˆØ�$‰;ØÐ'¨Ð,Ð-Ñ-ˆCØˆ
r   c                 óš   — t        j                  | fi |¤Ž}t        j                  j                  rt
        j                  | «      r|dz  }|S )Nr¸   )r   Úget_tensor_specializationr
   r	   Úuse_buffer_opsr{   r¶   )r´   Úkwargsr»   s      r   r½   z$HIPBackend.get_tensor_specializationº   s?   € ä×3Ñ3°CÑB¸6ÑBˆÜ�9‰9×#Ò#¬
×(@Ñ(@ÀÔ(EØ�3‰JˆCØˆ
r   c                 ó2  — t        j                  | j                  «      }|j                  «        t        j
                  j                  |«       t        j                  j                  |«       t        j                  |j                  «      s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)                  | d«       | S )NÚ	make_ttir)r   Úpass_managerÚcontextÚenable_debugr   ÚcommonÚadd_inlinerÚttirÚadd_rewrite_tensor_pointerr	   Úsupports_tdmr!   Ú(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   rÁ   zHIPBackend.make_ttirÁ   s÷   € ä�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�‰×!Ñ! "Ô%Ü�‰×.Ñ.¨rÔ2Ü×Ñ §¡Ô-Ü�K‰K×@Ñ@ÀÔDÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ Ô#Ü�‰×)Ñ)¨"Ô-Ü�‰×Ñ˜bÔ!Ü�‰×#Ñ# BÔ'Ü�‰×$Ñ$ RÔ(Ü�‰×#Ñ# BÔ'Ø
�‰ˆs�KÔ Øˆ
r   c                 óÐ  — t        j                  | j                  «      }|j                  «        t        j
                  j                  |d|j                  › �|j                  |j                  |j                  «       |j                  | d«       t        j                  | j                  «      }|j                  «        d}t        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                  j-                  ||j                  «       t"        j                  j                  j/                  |«       t"        j                  j                  j1                  |«       t        j                  j3                  |«       t        j4                  j7                  |«       t        j
                  j9                  |«       t        j4                  j7                  |«       t;        |j                  «      }t=        |j                  |«      }t"        j                  j                  j?                  ||j@                  «       t"        j                  j                  jC                  |||«       |r4t"        j                  j                  jE                  ||j                  «       t"        j                  j                  jG                  |«       t        j4                  j7                  |«       |jH                  jK                  «       dk7  rJ|jH                  jM                  d«      D ],  }t"        j                  j                  jO                  ||«       Œ. t        j                  j                  |«       t        j                  jQ                  |«       tS        |j                  «      rHt"        j                  j                  jU                  |«       t        j                  j                  |«       t"        j                  j                  jW                  |«       |rC|j@                  dkD  r4t"        j                  j                  jY                  ||j@                  «       tZ        j"                  j\                  r®t"        j                  j                  j_                  |«       t        j4                  j7                  |«       t"        j                  j                  ja                  ||j                  tZ        j"                  jb                  tZ        j"                  jd                  «       t"        j                  j                  jg                  |«       t"        j                  j                  ji                  |«       t        j4                  j7                  |«       t        j4                  jk                  |«       t        j4                  jm                  |«       |j                  | d«       | jo                  «       |d<   | S )	Nrˆ   Úmake_ttgir_earlyFrL   ú,r   Ú
make_ttgirÚtensordesc_meta)8r   rÂ   rÃ   rÄ   r   rÇ   Úadd_convert_to_ttgpuirr!   r/   rT   r3   rÒ   ÚttgpuirÚadd_coalesceÚadd_f32_dot_tcÚadd_remove_layout_conversionsÚadd_optimize_thread_localityr	   Úadd_accelerate_matmulrD   rE   Úadd_optimize_epilogueÚadd_optimize_dot_operandsÚadd_hoist_layout_conversionsÚadd_sink_layout_conversionsÚadd_fuse_nested_loopsrÅ   rË   rÏ   r*   r#   Úadd_schedule_loopsr2   Úadd_pipelineÚadd_coalesce_async_copyÚadd_convert_to_tensor_opsrM   ÚlowerÚsplitÚinsert_instruction_sched_hintsÚadd_reduce_data_duplicationr'   Úadd_in_thread_transposeÚadd_move_up_prologue_loadsÚadd_block_pingpongr
   r¾   Úadd_canonicalize_pointersÚadd_convert_to_buffer_opsÚuse_buffer_atomicsÚ%buffer_ops_analyze_small_tensor_rangeÚadd_fold_true_cmpiÚadd_prepare_if_combiningrÎ   rÐ   Úget_tensordesc_metadata)rÓ   r    rŠ   rÔ   ÚemuTF32r"   r    Úhints           r   rØ   zHIPBackend.make_ttgirÓ   s`  € ä�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�‰×*Ñ*¨2°°g·l±l°^Ð/DÀg×FWÑFWÐY`×YjÑYjØ+2×+;Ñ+;ô	=à
�‰ˆsÐ&Ô'Ü�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔØˆÜ�‰×#Ñ# BÔ'Ü�‰×%Ñ% b¨'Ô2Ü�‰×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Ü�
‰
×Ñ×4Ñ4°R¸¿¹ÔFÜ�
‰
×Ñ×7Ñ7¸Ô;Ü�
‰
×Ñ×6Ñ6°rÔ:ä�‰×,Ñ,¨RÔ0Ü�‰×'Ñ'¨Ô+Ü�‰×#Ñ# BÔ'Ü�‰×'Ñ'¨Ô+ä.¨w¯|©|Ó<ˆÜ9¸'¿,¹,ÈÓWÐä�
‰
×Ñ×-Ñ-¨b°'×2DÑ2DÔEÜ�
‰
×Ñ×'Ñ'¨¨NÐ<NÔOÙÜ�J‰J×Ñ×6Ñ6°r¸7¿<¹<ÔHÜ�
‰
×Ñ×4Ñ4°RÔ8Ü�‰×'Ñ'¨Ô+Ø× Ñ ×&Ñ&Ó(¨FÒ2Ø×-Ñ-×3Ñ3°CÓ8ò L�Ü—
‘
×"Ñ"×AÑAÀ"ÀdÕKðLä�‰×4Ñ4°RÔ8Ü�‰×2Ñ2°2Ô6Ü)¨'¯,©,Ô7Ü�J‰J×Ñ×6Ñ6°rÔ:Ü�N‰N×8Ñ8¸Ô<Ü�
‰
×Ñ×5Ñ5°bÔ9Ù '×"4Ñ"4°qÒ"8Ü�J‰J×Ñ×1Ñ1°"°g×6HÑ6HÔIä�9‰9×#Ò#Ü�J‰J×Ñ×8Ñ8¸Ô<Ü�M‰M×+Ñ+¨BÔ/Ü�J‰J×Ñ×8Ñ8ØØ—‘Ü—	‘	×,Ñ,Ü—	‘	×?Ñ?ô	ô 	�
‰
×Ñ×-Ñ-¨bÔ1Ü�
‰
×Ñ×3Ñ3°BÔ7Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ø
�‰ˆs�LÔ!Ø&)×&AÑ&AÓ&CˆÐ"Ñ#Øˆ
r   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                  j!                  |«       t        j                  j#                  |«       |j%                  |d«       |j'                  «       |d<   |S )NÚgluon_to_ttgirrÙ   )r   rÂ   rÃ   rÄ   r   ÚgluonrÆ   Úadd_resolve_auto_encodingsrÅ   Úadd_sccprÇ   Úadd_loop_aware_cserË   rÛ   Ú add_combine_tensor_select_and_ifr	   Úadd_warp_pipelineÚadd_allocate_warp_groupsrÒ   r÷   )Úsrcr    rŠ   rÓ   rÔ   s        r   rû   zHIPBackend.gluon_to_ttgir  sá   € àˆÜ�_‰_˜SŸ[™[Ó)ˆØ
�‰Ôä�‰× Ñ  Ô$Ü�‰×/Ñ/°Ô3Ü�‰×Ñ˜rÔ"Ü�‰×&Ñ& rÔ*Ü�‰×&Ñ& rÔ*Ü�‰×7Ñ7¸Ô;Ü�
‰
×Ñ×,Ñ,¨RÔ0Ü�‰×/Ñ/°Ô3à
�‰ˆsÐ$Ô%Ø&)×&AÑ&AÓ&CˆÐ"Ñ#Øˆ
r   c                 ó4  — | }t        j                  |j                  «      }|j                  «        t        j
                  j                  j                  ||j                  «       t        j
                  j                  j                  ||j                  «       t
        j                  j                  |«       t
        j                  j                  |«       t
        j                  j                  |«       t        j
                  j                  j                  |«       t         j"                  r+t         j"                  j%                  d||j                  «       d}t        j
                  j                  j'                  ||j                  |«       t        j
                  j                  j)                  ||j                  «       t
        j*                  j-                  |«       t
        j*                  j/                  |«       t
        j                  j1                  |«       t
        j                  j3                  |«       t
        j*                  j-                  |«       t
        j*                  j/                  |«       t
        j*                  j5                  |«       |j6                  j9                  «       dk7  r?t        j
                  j                  j;                  ||j                  |j<                  «       t         j"                  r+t         j"                  j%                  d||j                  «       t>        j@                  jB                  s9t>        j@                  jD                  st
        jF                  jI                  |«       t        j
                  j                  jK                  ||j                  |«       |jM                  |d«       t>        j@                  jD                  rÚt>        j@                  jB                  s`t        j                  |j                  «      }|j                  «        t
        jF                  jI                  |«       |jM                  |d«       t        j                  |j                  «      }|j                  «        t
        jF                  jO                  |«       |jM                  |d«       tQ        jR                  «        tQ        j                  «       }tQ        jT                  ||«      }t	        jV                  |«       d}t>        j@                  jX                  rd	}tQ        jZ                  |t        j\                  |j                  |«       t	        j^                  ||j                  «       t	        j`                  |d
«       t	        jb                  |dd«       t	        jb                  |dd«       t	        jb                  |dd«       t	        jb                  |d|jd                  dk(  «       |jg                  «       D �	cg c]  }	|	ji                  «       rŒ|	‘Œ }
}	|
d   jk                  t        jl                  «       |d   }|
d   jo                  d|› d�«       |jp                  }| js                  d«      }|�|}|
d   jo                  dd||jd                  z  › �«       d|j6                  ju                  d«      v r|
d   jo                  dd«       |
d   jo                  dd«       |
d   jo                  d|jv                  › d|jv                  › �«       |jx                  rd nd!}|
d   jo                  d"|«       t>        j@                  jX                  r'|
d   j{                  d	«       |
d   j}                  «        |j                  d#k7  rLt>        j                  j~                  s2t>        j                  j€                  st	        j‚                  |
d   «       t>        j@                  jX                  r\t…        t†        «      jˆ                  d$z  }t‹        |d%z  «      t‹        |d&z  «      t‹        |d'z  «      g}tQ        jŒ                  ||«       ne|jŽ                  rY|jŽ                  D ��cg c]  \  }}t	        j�                  ||«      sŒ|‘Œ  }}}t“        |«      dkD  rtQ        jŒ                  ||«       tQ        j”                  |tP        j–                  |j                  dg |j˜                  «       t	        jš                  |j                  «      r<|
d   j�                  d(«       |
d   j�                  d)«       |
d   j�                  d*«       t>        j                  jž                  rt	        j                   |
d   «       ||d+<   | js                  d,«      |d-<   | js                  d.«      xs d|d/<   | js                  d0«      xs d1|d2<   t	        j¢                  |«       t	        j¤                  |«       t‹        |«      S c c}	w c c}}w )3NÚttgpuir_to_llvmirTrL   Úllvmir_to_llvmÚ	make_llirzmake_llir.disable_line_infoz,make_llir.dump_ir_extract_di_local_variablesrJ   ú+xnackiô  Ú__oclc_finite_only_optFÚ__oclc_correctly_rounded_sqrt32Ú__oclc_unsafe_math_optÚ__oclc_wavefrontsize64rS   r   r3   zamdgpu-cluster-dimsz,1,1zttg.total-num-warpszamdgpu-flat-work-group-sizez1,zmemory-bound-attentionr×   zamdgpu-sched-strategyziterative-ilpzuniform-work-group-sizeÚtruezamdgpu-waves-per-euz, zpreserve-signr=   zdenormal-fp-math-f32Úgfx1250rU   z
asanrtl.bczocml.bczockl.bczamdgpu-no-workgroup-id-xzamdgpu-no-workgroup-id-yzamdgpu-no-workgroup-id-zr/   z
ttg.sharedrŸ   zttg.profile_scratch_memory_sizeÚprofile_scratch_sizez$ttg.profile_scratch_memory_alignmentr   Úprofile_scratch_align)Sr   rÂ   rÃ   rÄ   r	   r   rÛ   Úadd_update_async_wait_countr!   Úadd_warp_pipeline_conversionÚconvertÚadd_scf_to_cfrü   rÆ   Úadd_index_to_llvmirÚadd_allocate_shared_memoryr{   r«   ÚpatchÚadd_to_llvmirÚadd_warp_specialize_to_llvmrÅ   rË   rÎ   Úadd_cf_to_llvmirÚadd_arith_to_llvmirrÐ   rM   rê   Úlower_instruction_sched_hintsr2   r
   ÚcompilationÚdisable_line_infoÚ"dump_ir_extract_di_local_variablesÚllvmirÚadd_di_scopeÚadd_builtin_func_to_llvmirrÒ   Úadd_di_local_variabler   Úinit_targetsÚ	to_moduleÚattach_target_tripleÚenable_asanÚattach_datalayoutÚTARGET_TRIPLEÚset_isa_versionÚset_abi_versionÚset_bool_control_constantrT   Úget_functionsÚis_declarationÚset_calling_convÚCALLING_CONV_AMDGPU_KERNELÚadd_fn_attrr/   Úget_int_attrrë   r0   rF   Úadd_fn_target_featureÚadd_fn_asan_attrÚswap_mirÚdump_mirÚset_all_fn_arg_inregr   r]   r^   r`   Úlink_extern_libsr4   Úneed_extern_libÚlenÚoptimize_moduleÚOPTIMIZE_O3rB   Úhas_architected_sgprsÚremove_fn_attrÚscalarize_packed_fopsÚ#add_scalarize_packed_fops_llvm_passÚcleanup_bitcode_metadataÚdisable_print_inline)r  r    rŠ   rÓ   rÔ   Ú_HIPBackend__HIP_FTZrÃ   Úllvm_modÚtarget_featuresÚfnÚfnsÚcluster_dimÚtotal_warps_numÚtotal_num_warpsÚdenormal_modere   Úpathsrq   Úpaths                      r   r  zHIPBackend.make_llir(  so  € àˆä�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�
‰
×Ñ×6Ñ6°r¸7¿<¹<ÔHÜ�
‰
×Ñ×7Ñ7¸¸G¿L¹LÔIÜ�‰×$Ñ$ RÔ(Ü�‰× Ñ  Ô$Ü�‰×*Ñ*¨2Ô.ä�
‰
×Ñ×5Ñ5°bÔ9ä×%Ò%Ü×&Ñ&×,Ñ,Ð-@À"ÀcÇkÁkÔRð ˆ	Ü�
‰
×Ñ×(Ñ(¨¨W¯\©\¸9ÔEÜ�
‰
×Ñ×6Ñ6°r¸7¿<¹<ÔHÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!ä�‰×'Ñ'¨Ô+Ü�‰×*Ñ*¨2Ô.Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(à× Ñ ×&Ñ&Ó(¨FÒ2Ü�J‰J×Ñ×<Ñ<¸RÀÇÁÈw×OaÑOaÔbô ×%Ò%Ü×&Ñ&×,Ñ,Ð-=¸rÀ3Ç;Á;ÔOä× Ñ ×2Ò2¼5×;LÑ;L×;oÒ;oÜ�M‰M×&Ñ& rÔ*ä�
‰
×Ñ×5Ñ5°b¸'¿,¹,È	ÔRØ
�‰ˆs�KÔ ä×Ñ×?Ò?ä×$Ñ$×6Ò6Ü—_‘_ S§[¡[Ó1�Ø—‘Ô!Ü—‘×*Ñ*¨2Ô.Ø—‘�sÐ9Ô:ô —‘ §¡Ó-ˆBØ�O‰OÔÜ�M‰M×/Ñ/°Ô3Ø�F‰F�3ÐFÔGô 	×ÑÔÜ—,‘,“.ˆÜ—>‘> # 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‰×Ñ¤× >Ñ >Ô?Ø˜zÑ*ˆØˆA‰×ÑÐ0°[°MÀÐ2FÔGà!×+Ñ+ˆØ×*Ñ*Ð+@ÓAˆØÐ&Ø-ˆOØˆA‰×ÑÐ8¸B¸Èw×O`ÑO`Ñ?`Ð>aÐ:bÔcØ# w×'<Ñ'<×'BÑ'BÀ3Ó'GÑGØ�‰F×ÑÐ6¸ÔHØˆA‰×ÑÐ4°fÔ=ð 	ˆA‰×ÑÐ0°W×5IÑ5IÐ4JÈ"ÈW×MaÑMaÐLbÐ2cÔdØ+2×+EÒ+E™È6ˆØˆA‰×ÑÐ1°=ÔAÜ×Ñ×(Ò(Ø�‰F×(Ñ(¨Ô2Ø�‰F×#Ñ#Ô%ð �<‰<˜9Ò$¬e¯i©i×.@Ò.@ÄEÇIÁI×DVÒDVÜ×$Ñ$ S¨¡VÔ,ä×Ñ×(Ò(Ü!¤(›^×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Ü�5‹z˜AŠ~Ü×%Ñ% h°Ô6ä×Ñ˜X¤t×'7Ñ'7¸¿¹ÀrÈ2Èw×OgÑOgÔhô ×$Ñ$ W§\¡\Ô2Ø�‰F×!Ñ!Ð"<Ô=Ø�‰F×!Ñ!Ð"<Ô=Ø�‰F×!Ñ!Ð"<Ô=ä�9‰9×*Ò*Ü×3Ñ3°C¸±FÔ;ð !0ˆ�ÑØ ×-Ñ-¨lÓ;ˆ�ÑØ+.×+;Ñ+;Ð<]Ó+^Ò+cÐbcˆÐ'Ñ(Ø,/×,<Ñ,<Ð=cÓ,dÒ,iÐhiˆÐ(Ñ)ä×$Ñ$ XÔ.ô 	× Ñ  Ô*Ü�8‹}ÐùòW Qùó^ js   ØhØ,há>hâhc           	      óô  — t        j                  d| «      }t        |«      dk(  sJ ‚|d   |d<   g }d|j                  v rdnd}t	        j
                  | j                  d«      «      j                  «       }|d   d	z   |z   }t        j                  | t        j                  |j                  |||j                  |«      }t        j                  | t        j                  |j                  |||j                  |«       t        j                  j                  rwt        j                   t"        j$                  j'                  t        j                  j                  |d
z   «      t        j                  |j                  |||j                  d«      }	n=t        j(                  | t        j                  |j                  |||j                  d«      }	t        j                  j*                  rt-        d«       t-        |	«       |	S )Nz3define amdgpu_kernel void @([a-zA-Z_][a-zA-Z0-9_]*)r   r   rq   Úgfx11z-real-true16rJ   rj   rh   z.txtFz!// -----// AMDGCN Dump //----- //)ÚreÚfindallr:  r!   rm   rn   ro   rp   r   Útranslate_to_mirr	   r)  rB   Údump_sched_dagr
   r5  Útranslate_mir_to_asmÚosrM  rk   Útranslate_to_asmÚdump_amdgcnÚprint)
r  r    rŠ   ÚnamesÚflagsÚfeaturesÚir_hashÚdump_file_idrh   Úamdgcns
             r   Úmake_amdgcnzHIPBackend.make_amdgcnË  s–  € ô
 —
‘
ÐQÐSVÓWˆÜ�5‹z˜QŠÐˆØ  ™8ˆ�ÑàˆØ%,°·±Ñ%<‘>À"ˆÜ—.‘. §¡¨GÓ!4Ó5×?Ñ?ÓAˆØ˜Q‘x #‘~¨Ñ/ˆÜ×!Ñ! #¤s×'8Ñ'8¸'¿,¹,ÈÐRWÐY`×YqÑYqØ".ó0ˆä×Ñ˜C¤×!2Ñ!2°G·L±LÀ(ÈEÐSZ×SkÑSkØ(ô	*ä�9‰9×ÒÜ×.Ñ.¬r¯w©w¯|©|¼E¿I¹I×<NÑ<NØ<HÈ6Ñ<Qó0SÜTW×TeÑTeÐgn×gsÑgsØ/7¸À×@XÑ@XÐZ_óa‰Fô ×*Ñ*¨3´×0AÑ0AÀ7Ç<Á<ÐQYÐ[`Ø+2×+CÑ+CÀUóLˆFä�9‰9× Ò ÜÐ5Ô6Ü�&ŒMØˆr   c                 ó’  — d}t         j                  j                  rd}d|j                  v r|dz  }t	        j
                  | |j                  |«      }t        j                  «       5 }t        j                  «       5 }t        |j                  d«      5 }|j                  |«       d d d «       t	        j                  |j                  |j                  «       d d d «       t        |j                  d«      5 }|j                  «       }	d d d «       d d d «       	S # 1 sw Y   ŒtxY w# 1 sw Y   ŒNxY w# 1 sw Y   Œ+xY w# 1 sw Y   	S xY w)NrJ   r  rO  z,-real-true16ÚwbÚrb)r
   r  r'  r!   r	   Úassemble_amdgcnÚtempfileÚNamedTemporaryFileÚopenrq   ÚwriteÚ
link_hsacoÚread)
r  r    rŠ   rE  r�   Útmp_outÚtmp_inÚfd_inÚfd_outr»   s
             r   Ú
make_hsacozHIPBackend.make_hsacoè  s  € àˆÜ×Ñ×(Ò(Ø&ˆOØ�g—l‘lÑ"Ø˜Ñ.ˆOÜ×#Ñ# C¨¯©°ÓGˆÜ×(Ñ(Ó*ð 	$¨gÜ×,Ñ,Ó.ð :°&Ü˜&Ÿ+™+ tÓ,ð '°Ø—K‘K Ô&÷'ä—‘˜vŸ{™{¨G¯L©LÔ9÷:ô �g—l‘l DÓ)ð $¨VØ—k‘k“m�÷$÷	$ð ˆ
÷'ð 'ú÷:ð :ú÷$ð $ú÷	$ð ˆ
úsT   Á'D<Á<D$ÂDÂ%2D$ÃD<Ã5D0ÄD<ÄD!ÄD$Ä$D-	Ä)D<Ä0D9	Ä5D<Ä<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<   t        j                  j
                  �$t        j                  j                  ‰ |‰|d «       y y )Nc                 ó*   •— ‰j                  | |‰«      S r)   )rÁ   ©r  r    rŠ   rc   s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>û  s   ø€ °4·>±>À#ÀxÐQXÓ3Y€ r   rÇ   c                 ó*   •— ‰j                  | |‰«      S r)   )rØ   rq  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>ü  s   ø€ °D·O±OÀCÈÐSZÓ4[€ r   Úttgirc                 ó*   •— ‰j                  | |‰«      S r)   )rû   rq  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>þ  s   ø€ °D×4GÑ4GÈÈXÐW^Ó4_€ r   c                 ó*   •— ‰j                  | |‰«      S r)   )r  rq  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>ÿ  s   ø€ ¨t¯~©~¸cÀ8ÈWÓ/U€ r   Úllirc                 ó*   •— ‰j                  | |‰«      S r)   )r_  rq  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>   s   ø€ °×1AÑ1AÀ#ÀxÐQXÓ1Y€ r   r^  c                 ó*   •— ‰j                  | |‰«      S r)   )rn  rq  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>  s   ø€ °·±ÀÀXÈwÓ0W€ r   r�   )r   ÚTRITONÚGLUONr
   rŽ   Úadd_stages_inspection_hook)rc   ÚstagesrŠ   r–   s   ` ` r   Ú
add_stageszHIPBackend.add_stagesù  sƒ   ù€ Ø”x—‘Ò&ÜYˆF�6‰NÜ[ˆF�7ŠOØœŸ™Ò'Ü_ˆF�7‰OÜUˆˆv‰ÜYˆˆxÑÜWˆˆw‰Ü�=‰=×3Ñ3Ð?Ü�M‰M×4Ñ4°T¸6À7ÈHÐVZÕ[ð @r   c                 ó   — | j                   › S r)   r   )rc   s    r   rt   zHIPBackend.hash  s   € à—+‘+�Ðr   )!ru   rv   rw   r«   Ú%supports_native_tensor_specializationÚstaticmethodr   r~   rƒ   r`   r‹   r   r�   r¡   r¤   r   r   r¨   rª   r¶   r¹   r½   rÁ   rØ   rû   r  r_  rn  r}  Ú	functoolsÚ	lru_cachert   Ú__classcell__)r†   s   @r   r{   r{   k   sS  ø„ Ø€OØ,1Ð)àð' 	ò 'ó ð'ð"˜yð "¨Tõ "ð
%¨#ó %ð" Só "ò2
ò?ð>  S¨* _Ñ 5ó >ò
:ð
 ñó ðð ñó ðð ñó ðð ñó ðð" ñ?ó ð?ðB ñó ðð$ ñ`ó ð`ðD ñó ðð8 ñó ðò 
\ð €Y×ÑÓñ ó ô r   r{   )!Útriton.backends.compilerr   r   r   Útriton._C.libtritonr   r   r   r	   Útritonr
   Údataclassesr   Útypingr   r   r   Útypesr   rU  rm   rd  rP  r�  r[   Úpathlibr   r   r#   r'   r*   r-   r{   r   r   r   ú<module>r‹     s‚   ðß EÑ Eß 5Ó 5Ý Ý !ß #Ñ #Ý Û 	Û Û Û 	Û Û Ý ð0˜Yó 0òRò
rò$ñ �$Ô÷D?ð D?ó ðD?ôN\ �õ \ r   