Ë
    5^(h¿J  ã                   óè   — d dl 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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 e	d¬«       G d„ d«      «       Z G d„ de«      Zy)é    )ÚBaseBackendÚ	GPUTarget)ÚirÚpassesÚllvmÚamd)Ú	dataclass)ÚAnyÚDictÚTuple)Ú
ModuleTypeN)ÚPathÚtargetc                 ó   — d„ S )Nc                  ó   — y)N©é   r   r   © )ÚlhsTypeÚrhsTypes     úZ/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/triton/backends/amd/compiler.pyú<lambda>zmin_dot_size.<locals>.<lambda>   s   � ó    r   ©r   s    r   Úmin_dot_sizer      s   € á-Ð-r   c                 óF   — | dk(  rdnd}t        j                  d|«      dk(  S )NÚgfx942Ú1Ú0ÚTRITON_HIP_USE_BLOCK_PINGPONG)ÚosÚgetenv)ÚarchÚdefaults     r   Úis_pingpong_enabledr%      s'   € Ø˜XÒ%‰c¨3€GÜ�9‰9Ð4°gÓ>À#ÑEÐEr   T)Úfrozenc                   óŠ  — e Zd ZU dZeed<   dZeed<   dZeed<   dZeed<   dZ	ee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_ctasr   Únum_buffers_warp_specÚnum_consumer_groupsÚreg_dec_producerÚreg_inc_consumerNÚextern_libsr   Úcluster_dimsFÚdebugTÚsanitize_overflowr#   )Úfp8e5Úsupported_fp8_dtypesr   Údeprecated_fp8_dtypesÚieeeÚdefault_dot_input_precision)r:   Ú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ÚnoneÚinstruction_sched_variantc                 óx  — t        t        «      j                  dz  }| j                  €i nt	        | j                  «      }d| j
                  v sd| j
                  v sd| j
                  v rdnd}t        j                  | d|«       | j
                  dk(  rd	n| j                  }t        j                  | d
|«       ddg}|D ]  }t        ||› d�z  «      ||<   Œ t        j                  | dt        |j                  «       «      «       | j                  dkD  r| j                  | j                  d	z
  z  dk(  sJ d«       ‚y )NÚlibÚgfx10Úgfx11Úgfx12é    é@   Ú	warp_sizeÚgfx950r   r@   ÚocmlÚocklz.bcr3   r   znum_warps must be a power of 2)r   Ú__file__Úparentr3   Údictr#   ÚobjectÚ__setattr__r@   ÚstrÚtupleÚitemsr*   )ÚselfÚdefault_libdirr3   rN   r@   ÚlibsrH   s          r   Ú__post_init__zHIPOptions.__post_init__G   s'  € Üœh›×.Ñ.°Ñ6ˆØ ×,Ñ,Ð4‘b¼$¸t×?OÑ?OÓ:Pˆà! T§Y¡YÑ.°'¸T¿Y¹YÑ2FÈ'ÐUY×U^ÑU^ÑJ^‘BÐdfˆ	Ü×Ñ˜4 ¨iÔ8à—Y‘Y (Ò*‘°·
±
ˆÜ×Ñ˜4 ¨%Ô0Ø˜ÐˆØò 	AˆCÜ" >°s°e¸3°KÑ#?Ó@ˆK˜Òð	Aä×Ñ˜4 ´°k×6GÑ6GÓ6IÓ0JÔKØ�~‰~ Ò! t§~¡~¸¿¹È!Ñ9KÑ'LÐQRÒ&Rð 	0Ø/ó	0ÐRÐ&Rr   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__rY   ÚhashlibÚsha256ÚencodeÚ	hexdigest)rZ   ÚnameÚvalÚkeys       r   ÚhashzHIPOptions.hashW   s]   € Ø�h‰h¸¿¹×9LÑ9LÓ9N×O©I¨D°#˜4˜&  # šÓOÓPˆÜ�~‰~˜cŸj™j¨Ó1Ó2×<Ñ<Ó>Ð>ùó Ps   ©A4
)%Ú__name__Ú
__module__Ú__qualname__r*   ÚintÚ__annotations__r+   r-   r.   r/   r0   r1   r2   r3   rT   r4   rX   r5   Úboolr6   r#   rW   r8   r   r9   r;   r<   r=   r>   r?   r@   rA   rB   rD   rF   r]   rk   r   r   r   r(   r(      s"  … à€IˆsÓØ€L�#ÓØ€J�ÓØ€HˆcÓØ!"Ð˜3Ó"Ø Ð˜Ó ØÐ�cÓØÐ�cÓØ€K�ÓØ#€L�%Ó#Ø€Eˆ4ÓØ"Ð�tÓ"Ø€Dˆ#ÓØ'2Ð˜% ™*Ó2Ø(*Ð˜5 ™:Ó*Ø'-Ð Ó-Ø/9Ð  %¨¡*Ó9Ø!Ð�dÓ!Ø$)Ð˜TÓ)Ø !Ð˜#Ó!Ø€Eˆ3ƒNØ$Ð˜Ó$Ø)*Ð! 3Ó*Ø€L�#Óð& &,Ð˜sÓ+ò0ó ?r   r(   c                   ón  ‡ — e Zd Zedefd„«       Zdeddfˆ fd„Zdefd„Zd„ Z	d„ Z
deeef   fd	„Zd
„ Ze ej"                  «       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 )NrC   )Ú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#   rW   Ú
binary_ext)rZ   r   Ú	__class__s     €r   r{   zHIPBackend.__init__b   s+   ø€ Ü‰Ñ˜Ô Ü˜&Ÿ+™+¤sÔ+Ð+Ð+Ø!ˆ�r   c                 ó  — dt        j                  d| j                  j                  «      i}| j                  j                  dv rBt	        t
        j                  «      }|j                  dh«       t        t        |«      «      |d<   d|vr‡t	        t
        j                  «      }| j                  j                  dv r|j                  h d£«       n+| j                  j                  dv r|j                  d	d
h«       t        t        |«      «      |d<   d|vrt        j                  dd«      dk(  |d<   |j                  t
        j                  j                  «       D �ci c]  }||v sŒ||   €Œ|||   “Œ c}«       t        di |¤ŽS c c}w )Nr#   ÚTRITON_OVERRIDE_ARCH)Úgfx940Úgfx941r   Útf32r<   r8   >   Úfp8e4b8Úfp8e4nvÚfp8e5b16rO   r…   r7   r=   ÚTRITON_DEFAULT_FP_FUSIONr   r   )r!   r"   r   r#   Úsetr(   r<   ÚupdaterX   Úsortedr8   Ú__dataclass_fields__Úkeys)rZ   ÚoptsÚargsr<   r8   Úks         r   Úparse_optionszHIPBackend.parse_optionsg   s`  € ØœŸ	™	Ð"8¸$¿+¹+×:JÑ:JÓKÐLˆð �;‰;×ÑÐ=Ñ=Ü+.¬z×/VÑ/VÓ+WÐ(Ø(×/Ñ/°°Ô9Ü38¼Ð@\Ó9]Ó3^ˆDÐ/Ñ0à!¨Ñ-Ü#&¤z×'FÑ'FÓ#GÐ Ø�{‰{×ÑÐ#AÑAØ$×+Ñ+Ò,NÕOØ—‘×!Ñ! hÑ/Ø$×+Ñ+¨Y¸Ð,@ÔAÜ+0´Ð8LÓ1MÓ+NˆDÐ'Ñ(à TÑ)Ü')§y¡yÐ1KÈSÓ'QÐUXÑ'XˆDÐ#Ñ$Ø�‰¬×)HÑ)H×)MÑ)MÓ)OÖu AÐSTÐX\ÒS\ÐaeÐfgÑahÑat�Q˜˜Q™‘ZÒuÔvÜÑ!˜DÑ!Ð!ùò vs   Å	FÅ'FÅ-Fc                 óœ   — |j                   |j                  |j                  |j                  d   |j                  d   |j                  d   fS )Nr   r   r,   )r*   r.   Úsharedr4   )rZ   Úmetadatas     r   Úpack_metadatazHIPBackend.pack_metadata}   sO   € à×ÑØ×ÑØ�O‰OØ×!Ñ! !Ñ$Ø×!Ñ! !Ñ$Ø×!Ñ! !Ñ$ð
ð 	
r   c                 ó4   — dt        | j                  «      i}|S )Nr   )r   r   )rZ   ÚoptionsÚcodegen_fnss      r   Úget_codegen_implementationz%HIPBackend.get_codegen_implementation‡   s   € Ø%¤|°D·K±KÓ'@ÐAˆØÐr   c                 ó   — ddl m} d|iS )Nr   )Ú	libdeviceztriton.language.extra.libdevice)Útriton.language.extra.hiprš   )rZ   rš   s     r   Úget_module_mapzHIPBackend.get_module_map‹   s   € Ý7à1°9Ð=Ð=r   c                 ó.   — t        j                  |«       y ©N)r   Úload_dialects)rZ   Úctxs     r   rŸ   zHIPBackend.load_dialects�   s   € Ü×Ñ˜#Õr   c                  óH   — t         j                  j                  dd«      dk(  S )NÚAMDGCN_USE_BUFFER_OPSr   r   )r!   ÚenvironÚgetr   r   r   Úuse_buffer_opszHIPBackend.use_buffer_ops“   s   € ô �z‰z�~‰~Ð5°sÓ;¸sÑBÐBr   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_rangerL   )r   Ú
parse_attr)ÚdescÚrets     r   r²   zHIPBackend.parse_attr£   s1   € ä×$Ñ$ TÓ*ˆØ�$‰;ØÐ'¨Ð,Ð-Ñ-ˆCØˆ
r   c                 óš   — t        j                  | |fi |¤Ž}t        j                  «       r|dk(  rt        j	                  | «      r|dz  }|S )NÚtensorr±   )r   Úget_arg_specializationrs   r¥   r¯   )r­   ÚtyÚkwargsr´   s       r   r·   z!HIPBackend.get_arg_specializationª   sI   € ä×0Ñ0°°bÑC¸FÑCˆô ×$Ñ$Ô&¨2°ª>¼j×>VÑ>VÐWZÔ>[Ø�3‰JˆCØˆ
r   c                  óP  — t        j                  d«      } | �t        | «      }|j                  «       r|S t        t        «      j
                  dz  }|j                  «       r|S t        d«      }|j                  «       r|S t        d«      }|j                  «       r|S t        d«      ‚)NÚTRITON_HIP_LLD_PATHz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"   r   Úis_filerR   rS   Ú	Exception)Úlld_env_pathÚllds     r   Úpath_to_rocm_lldzHIPBackend.path_to_rocm_lld³   s’   € ô —y‘yÐ!6Ó7ˆØÐ#Ü�|Ó$ˆ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                  |«       |j!                  | «       | S rž   )r   Úpass_managerÚcontextÚenable_debugr   ÚcommonÚadd_inlinerÚttirÚadd_rewrite_tensor_pointerÚadd_canonicalizerÚadd_combineÚadd_reorder_broadcastÚadd_cseÚadd_licmÚadd_symbol_dceÚadd_loop_unrollÚrun)Úmodr“   r–   Úpms       r   Ú	make_ttirzHIPBackend.make_ttirÇ   sÐ   € ä�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�‰×!Ñ! "Ô%Ü�‰×.Ñ.¨rÔ2Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ Ô#Ü�‰×)Ñ)¨"Ô-Ü�‰×Ñ˜bÔ!Ü�‰×Ñ˜rÔ"Ü�‰×$Ñ$ 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/        t1        j2                  dd«      «      }t/        t1        j2                  dd«      «      }|j4                  dk(  rdx}}t!        j6                  |j                  «      rk|j8                  dk7  sJ d	«       ‚t         j                  j                  j;                  ||j8                  ||«       t        j<                  j?                  |«       |j4                  jA                  «       d
k7  r4t         j                  j                  jC                  ||j4                  «       t        j                  j+                  |d«       t        j                  j                  |«       t        j                  jE                  |«       t!        j6                  |j                  «      rxt         j                  j                  jG                  |«       tI        |j                  «      }|r8|j8                  dk(  r)t         j                  j                  jK                  |«       tL        jO                  «       r|t         j                  j                  jQ                  |«       t        j<                  j?                  |«       t         j                  j                  jS                  ||j                  «       t        j<                  j?                  |«       t        j<                  jU                  |«       t        j<                  jW                  |«       |j                  | «       | S )Nzhip:TÚTRITON_HIP_GLOBAL_PREFETCHr   ÚTRITON_HIP_LOCAL_PREFETCHzlocal-prefetchr   r   zÕTriton AMD backend pipeliner has been updated. We used to trigger software pipelining with num_stages == 0. Now it will not happen anymore; please update to use num_stages == 2 for equivalent behavior in the past.rE   r,   ),r   rÂ   rÃ   rÄ   r   rÇ   Úadd_convert_to_ttgpuirr#   r*   rN   r.   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_conversionsro   r!   r"   rF   Úhas_matrix_core_featurer-   Úadd_stream_pipelinerÅ   rÉ   ÚlowerÚinsert_instruction_sched_hintsÚadd_reduce_data_duplicationÚadd_reorder_instructionsr%   Úadd_block_pingpongrs   r¥   Úadd_canonicalize_pointersÚadd_convert_to_buffer_opsrÌ   rÎ   )rÑ   r“   r–   rÒ   Úglobal_prefetchÚlocal_prefetchÚuse_block_pingpongs          r   Ú
make_ttgirzHIPBackend.make_ttgir×   sZ  € ä�_‰_˜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¸Ô;äœbŸi™iÐ(DÀcÓJÓKˆÜœRŸY™YÐ'BÀCÓHÓIˆð ×,Ñ,Ð0@Ò@Ø/0Ð0ˆO˜nä×&Ñ& w§|¡|Ô4Ø×%Ñ%¨Ò*ð Qð .Pó QÐ*ô
 �J‰J×Ñ×2Ñ2°2°w×7IÑ7IÈ?Ð\jÔkÜ�M‰M×+Ñ+¨BÔ/Ø×,Ñ,×2Ñ2Ó4¸Ò>Ü�J‰J×Ñ×=Ñ=¸bÀ'×BcÑBcÔdÜ�‰×0Ñ0°°TÔ:Ü�‰×4Ñ4°RÔ8Ü�‰×2Ñ2°2Ô6Ü×&Ñ& w§|¡|Ô4Ü�J‰J×Ñ×7Ñ7¸Ô;Ü!4°W·\±\Ó!BÐÙ! g×&8Ñ&8¸AÒ&=Ü—
‘
×"Ñ"×5Ñ5°bÔ9ä×$Ñ$Ô&Ü�J‰J×Ñ×8Ñ8¸Ô<Ü�M‰M×+Ñ+¨BÔ/Ü�J‰J×Ñ×8Ñ8¸¸W¿\¹\ÔJÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ø
�‰ˆsŒØˆ
r   c                 óp  — | }t        j                  |j                  «      }|j                  «        t        j
                  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                  |j0                  «       t2        j4                  j7                  dd«      dk(  rt
        j8                  j;                  |«       t        j
                  j                  j=                  ||«       |j?                  |«       tA        jB                  «        tA        j                  «       }tA        jD                  ||«      }t	        jF                  |«       d}	t2        j4                  j7                  dd«      dk(  rd	}	tA        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|«       t2        j4                  j7                  dd«      dk(  r'|d   je                  d	«       |d   jg                  «        t	        jh                  |d   «       t2        j4                  j7                  dd«      dk(  r\tk        tl        «      jn                  dz  }tq        |dz  «      tq        |dz  «      tq        |dz  «      g}tA        jr                  ||«       nW|jt                  rK|jt                  D ��cg c]  \  }}t	        jv                  ||«      sŒ|‘Œ  }}}tA        jr                  ||«       tA        jx                  |t@        jz                  |j                  dg |j|                  «       | j                  d«      |d<   t	        j€                  |«       t	        j‚                  |«       tq        |«      S c c}
w c c}}w )Nr   TrE   ÚTRITON_DISABLE_LINE_INFOr   Ú ÚTRITON_ENABLE_ASANr   ú+xnackiô  Ú__oclc_finite_only_optFÚ__oclc_correctly_rounded_sqrt32Ú__oclc_unsafe_math_optÚ__oclc_wavefrontsize64rM   zamdgpu-flat-work-group-sizez1,zamdgpu-waves-per-euzpreserve-signr:   zdenormal-fp-math-f32rH   z
asanrtl.bczocml.bczockl.bcz
ttg.sharedr’   )Br   rÂ   rÃ   rÄ   r   r   rØ   Ú%add_decompose_unsupported_conversionsr#   Úadd_optimize_lds_usageÚ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Î   rF   râ   Úlower_instruction_sched_hintsr-   r!   r£   r¤   ÚllvmirÚadd_di_scopeÚadd_builtin_func_to_llvmirrÐ   r   Úinit_targetsÚ	to_moduleÚattach_target_tripleÚattach_datalayoutÚTARGET_TRIPLEÚset_isa_versionÚset_abi_versionÚset_bool_control_constantrN   Úget_functionsÚis_declarationÚset_calling_convÚCALLING_CONV_AMDGPU_KERNELÚadd_fn_attrr*   r+   rA   Úadd_fn_target_featureÚadd_fn_asan_attrÚset_all_fn_arg_inregr   rR   rS   rW   Úlink_extern_libsr3   Úneed_extern_libÚoptimize_moduleÚOPTIMIZE_O3r=   Úget_int_attrÚcleanup_bitcode_metadataÚdisable_print_inline)Úsrcr“   r–   rÑ   rÒ   Úcustom_lds_sizeÚ_HIPBackend__HIP_FTZrÃ   Úllvm_modÚtarget_featuresÚfnÚfnsÚdenormal_moder[   Úpathsrh   Úpaths                    r   Ú	make_llirzHIPBackend.make_llir  sï  € àˆä�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�
‰
×Ñ×@Ñ@ÀÀWÇ\Á\ÔRð ˆÜ�
‰
×Ñ×1Ñ1°"°g·l±lÀOÔTÜ�‰×$Ñ$ RÔ(Ü�‰×*Ñ*¨2Ô.ä�‰×1Ñ1°"Ô5ð ˆ	Ü�
‰
×Ñ×(Ñ(¨¨W¯\©\¸9ÔEÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!ä�‰×'Ñ'¨Ô+Ü�‰×*Ñ*¨2Ô.Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ø×,Ñ,×2Ñ2Ó4¸Ò>Ü�J‰J×Ñ×<Ñ<¸RÀÇÁÈw×OaÑOaÔbÜ�:‰:�>‰>Ð4°cÓ:¸cÒAÜ�M‰M×&Ñ& rÔ*Ü�
‰
×Ñ×5Ñ5°b¸)ÔDØ
�‰ˆsŒô 	×ÑÔÜ—,‘,“.ˆÜ—>‘> # wÓ/ˆÜ× Ñ  Ô*ØˆÜ�:‰:�>‰>Ð.°Ó4¸Ò;Ø&ˆ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Ü�:‰:�>‰>Ð.°Ó4¸Ò;Ø�‰F×(Ñ(¨Ô2Ø�‰F×#Ñ#Ô%ô
 	× Ñ   Q¡Ô(ä�:‰:�>‰>Ð.°Ó4¸Ò;Ü!¤(›^×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ð !×-Ñ-¨lÓ;ˆ�Ñä×$Ñ$ XÔ.ô 	× Ñ  Ô*Ü�8‹}ÐùòY Qùó@ js   Ð>Z-ÑZ-×,Z2ØZ2c           	      óN  — t        j                  d| «      }t        |«      dk(  sJ ‚|d   |d<   t        j                  | t
        j                  |j                  dg |j                  d«      }t        j                  j                  dd«      d	k(  rt        d
«       t        |«       |S )Nz3define amdgpu_kernel void @([a-zA-Z_][a-zA-Z0-9_]*)r   r   rh   rï   FÚAMDGCN_ENABLE_DUMPr   r   z!// -----// AMDGCN Dump //----- //)ÚreÚfindallÚlenr   Útranslate_to_asmr   r  r#   r=   r!   r£   r¤   Úprint)r  r“   r–   ÚnamesÚamdgcns        r   Úmake_amdgcnzHIPBackend.make_amdgcny  s�   € ô
 —
‘
ÐQÐSVÓWˆÜ�5‹z˜QŠÐˆØ  ™8ˆ�Ñä×&Ñ& s¬C×,=Ñ,=¸w¿|¹|ÈRÐQSÐU\×UmÑUmÐotÓuˆÜ�:‰:�>‰>Ð.°Ó4¸Ò;ÜÐ5Ô6Ü�&ŒMØˆr   c                 ó²  — d}t         j                  j                  dd«      dk(  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ð   r   r   rñ   Úwbz-flavorÚgnuz-sharedz-oÚrb)r!   r£   r¤   r   Úassemble_amdgcnr#   rs   rÀ   ÚtempfileÚNamedTemporaryFileÚopenrh   ÚwriteÚ
subprocessÚ
check_callÚread)r  r“   r–   r  ry   Ú	rocm_pathÚtmp_outÚtmp_inÚfd_inÚfd_outr´   s              r   Ú
make_hsacozHIPBackend.make_hsacoˆ  s)  € àˆÜ�:‰:�>‰>Ð.°Ó4¸Ò;Ø&ˆ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   Á1EÂD4ÂD(Â/8D4Ã'EÄE ÄEÄ(D1Ä-D4Ä4D=	Ä9EÅ E		ÅEÅEc                 ób   ‡ ‡— ˆˆ fd„|d<   ˆˆ fd„|d<   ˆˆ fd„|d<   ˆˆ fd„|d<   ˆˆ fd	„|d
<   y )Nc                 ó*   •— ‰j                  | |‰«      S rž   )rÓ   ©r  r“   r–   rZ   s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>š  ó   ø€ ¨t¯~©~¸cÀ8ÈWÓ/U€ r   rÇ   c                 ó*   •— ‰j                  | |‰«      S rž   )rì   rC  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>›  ó   ø€ °·±ÀÀXÈwÓ0W€ r   Úttgirc                 ó*   •— ‰j                  | |‰«      S rž   )r$  rC  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>œ  rD  r   Úllirc                 ó*   •— ‰j                  | |‰«      S rž   )r.  rC  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>�  s   ø€ °×1AÑ1AÀ#ÀxÐQXÓ1Y€ r   r-  c                 ó*   •— ‰j                  | |‰«      S rž   )r@  rC  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>ž  rF  r   ry   r   )rZ   Ústagesr–   s   ` `r   Ú
add_stageszHIPBackend.add_stages™  s1   ù€ ÜUˆˆv‰ÜWˆˆw‰ÜUˆˆv‰ÜYˆˆxÑÜWˆˆwŠr   c                 óz   — t        j                  t        j                  «       dgd¬«      }|› d| j                  › �S )Nz	--versionra   )Úencodingr`   )r8  Úcheck_outputrs   rÀ   r   )rZ   Úversions     r   rk   zHIPBackend.hash   s8   € ä×)Ñ)¬:×+FÑ+FÓ+HÈ+Ð*VÐahÔiˆØ�˜!˜DŸK™K˜=Ð)Ð)r   )rl   rm   rn   Ústaticmethodr   rv   r{   r
   r�   r”   r˜   r   rW   r   rœ   rŸ   Ú	functoolsÚ	lru_cacher¥   r¯   r²   r·   rÀ   rÓ   rì   r$  r.  r@  rM  rk   Ú__classcell__)r~   s   @r   rs   rs   \   sg  ø„ àð' 	ò 'ó ð'ð"˜yð "¨Tõ "ð
" Só "ò,
òð>  S¨* _Ñ 5ó >ò
ð Ø€Y×ÑÓñCó ó ðCð ñó ðð ñó ðð ñó ðð ñsó ðsð& ñó ðð ñ3ó ð3ðj ñió ðiðV ñó ðð ñó ðò Xð €Y×ÑÓñ*ó ô*r   rs   )Útriton.backends.compilerr   r   Útriton._C.libtritonr   r   r   r   Údataclassesr	   Útypingr
   r   r   Útypesr   rd   r4  r!   r'  r8  rS  Úpathlibr   r   r%   r(   rs   r   r   r   ú<module>r\     sn   ðß ;ß 5Ó 5Ý !ß #Ñ #Ý Û Û Û 	Û 	Û Û Ý ð.˜ó .ò
Fñ
 �$Ô÷??ð ??ó ð??ôDG*�õ G*r   