§
    NŠtj~^  ã                   ó  — 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 dS )é    )ÚBaseBackendÚ	GPUTargetÚLanguage)ÚirÚpassesÚllvmÚamd)Úknobs)Ú	dataclass)ÚAnyÚDictÚTuple)Ú
ModuleTypeN)ÚPathÚtargetc                 ó   — d„ S )Nc                 ó   — dS )N)é   r   r   © )Úlhs_typeÚrhs_types     úZ/var/www/html/CA-Chatbot/venv/lib/python3.11/site-packages/triton/backends/amd/compiler.pyú<lambda>z"get_min_dot_size.<locals>.<lambda>   s   €  i€ ó    r   ©r   s    r   Úget_min_dot_sizer      s   € ð 0Ð/Ð/r   c                 óf   — t           j        j        €| dk    p	| dk    o|du nt           j        j        S )NÚgfx942Úgfx950T)r
   r	   Úuse_block_pingpong)ÚarchÚuse_async_copys     r   Úis_pingpong_schedule_enabledr#      s:   € åŒ9Ô'Ð/ð �HÒÐM ¨Ò!1Ð!L°nÈÐ6LøÝ5:´YÔ5QðRr   c                 óR   — t           j        j        €| dk    nt           j        j        S )Nr   )r
   r	   Úuse_in_thread_transpose©r!   s    r   Úis_in_thread_transpose_enabledr'      s#   € Ý!&¤Ô!BÐ!JˆD�HÒÐÕPUÔPYÔPqÐqr   c                 ó$   — t           j        j        S ©N)r
   r	   r"   r&   s    r   Úis_async_copy_enabledr*      s   € ÝŒ9Ô#Ð#r   T)Úfrozenc                   óf  — 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d	S )&Ú
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                               | d|¦  «         | j        dk    r| j        | j        dz
  z  dk    s
J d	¦   «         ‚| j        d
k    rD| j        dk    r9t          j        d| j        › d�¦  «         t                               | dd¦  «         t          t          ¦  «        j
        dz  }| j        €i nt          | j        ¦  «        }dD ]}t          ||› d�z  ¦  «        ||<   Œt                               | dt          |                     ¦   «         ¦  «        ¦  «         d S )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   su  € Ý˜œ	 ! B $œÑ(Ô(ˆ	Ø# rš/˜/�B�B¨rˆ	Ý×Ò˜4 ¨iÑ8Ô8Ð8ØŒ~ Ò!Ð! t¤~¸¼È!Ñ9KÑ'LÐQRÒ&RÐ&RÐ&RØ,ñ 'SÔ&RÐRð ŒI˜Ò!Ð!¨¬
°aª¨ÝŒMð zÐvzô  wAð  zð  zð  zñô ð õ ×Ò˜t W¨aÑ0Ô0Ð0å�h™œÔ.°Ñ6ˆØ Ô,Ð4�b�b½$¸tÔ?OÑ:PÔ:PˆØ#ð 	Að 	AˆCÝ" >°s°K°K°KÑ#?Ñ@Ô@ˆK˜ÑÐÝ×Ò˜4 µ°k×6GÒ6GÑ6IÔ6IÑ0JÔ0JÑKÔKÐKÐKÐKr   c                 óà   — d                      d„ | j                             ¦   «         D ¦   «         ¦  «        }t          j        |                     d¦  «        ¦  «                             ¦   «         S )NÚ_c                 ó"   — g | ]\  }}|› d |› �‘ŒS )ú-r   )Ú.0ÚnameÚvals      r   ú
<listcomp>z#HIPOptions.hash.<locals>.<listcomp>g   s&   € ÐOÐOÐO©I¨D°#˜4˜˜ #˜˜ÐOÐOÐOr   úutf-8)ÚjoinÚ__dict__rb   ÚhashlibÚsha256ÚencodeÚ	hexdigest)rc   Úkeys     r   ÚhashzHIPOptions.hashf   sX   € Ø�hŠhÐOÐO¸¼×9LÒ9LÑ9NÔ9NÐOÑOÔOÑPÔPˆÝŒ~˜cŸjšj¨Ñ1Ô1Ñ2Ô2×<Ò<Ñ>Ô>Ð>r   ) Ú__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   rw   r   r   r   r-   r-   #   s   € € € € € € à€IˆsÐÐÑØ€L�#ÐÐÑØ€J�ÐÐÑØ€HˆcÐÐÑØ€K�ÐÐÑØ€Eˆ4ÐÐÑØ"Ð�tÐ"Ð"Ñ"Ø€Dˆ#ÐÐÑð
 (SÐ˜% œ*ÐRÐRÑRØ46Ð% u¨S¤zÐ6Ð6Ñ6Ø'-Ð Ð-Ð-Ñ-Ø/KÐ  %¨¤*ÐKÐKÑKØ!Ð�dÐ!Ð!Ñ!Ø$)Ð˜TÐ)Ð)Ñ)Ø !Ð˜#Ð!Ð!Ñ!Ø€Eˆ3€N€N�NØ$Ð˜Ð$Ð$Ñ$Ø)*Ð! 3Ð*Ð*Ñ*Ø€L�#ÐÐÑØ "Ð˜#Ð"Ð"Ñ"ð*  €M�3ÐÐÑðLð Lð Lð&?ð ?ð ?ð ?ð ?r   r-   c                   óŽ  ‡ — 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                 ó�   •— t          ¦   «                              |¦  «         t          |j        t          ¦  «        sJ ‚d| _        d S )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Ð$Ð$Ð$r   c                 ó€  ‡— dt           j        j        p| j        j        i}‰                     dd¦  «        dk    r:t          j        | j        j        ¦  «        st          d| j        j        › �¦  «        ‚| j        j        dk    rNt          t          j        ¦  «        }|                     dh¦  «         t          t          |¦  «        ¦  «        |d<   d‰vr)t          t          t          j        ¦  «        ¦  «        |d<   | j        j        d	k    rOt          t          j        ¦  «        }|                     d
dh¦  «         t          t          |¦  «        ¦  «        |d<   d‰vrt           j        j        |d<   |                     ˆfd„t          j                             ¦   «         D ¦   «         ¦  «         t          di |¤ŽS )Nr!   r3   r   znum_ctas > 1 not supported on r   Útf32rA   r;   r   r9   r:   r<   rB   c                 ó:   •— i | ]}|‰v ¯‰|         ®|‰|         “ŒS r)   r   )rk   ÚkÚoptss     €r   ú
<dictcomp>z,HIPBackend.parse_options.<locals>.<dictcomp>‘   s7   ø€ ÐuÐuÐu AÐSTÐX\ÐS\ÐS\ÐaeÐfgÔahÐat�Q˜˜QœÐatÐatÐatr   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   r“   ÚargsrA   r<   s    `   r   Úparse_optionszHIPBackend.parse_options{   sŸ  ø€ Ø�œÔ3ÐG°t´{Ô7GÐHˆà�8Š8�J Ñ"Ô" QÒ&Ð&­sÔ/LÈTÌ[ÔM]Ñ/^Ô/^Ð&ÝÐP¸d¼kÔ>NÐPÐPÑQÔQÐQð Œ;Ô˜xÒ'Ð'Ý+.­zÔ/VÑ+WÔ+WÐ(Ø(×/Ò/°°Ñ9Ô9Ð9Ý38½Ð@\Ñ9]Ô9]Ñ3^Ô3^ˆDÐ/Ñ0à!¨Ð-Ð-Ý+0µ½
Ô8WÑ1XÔ1XÑ+YÔ+YˆDÐ'Ñ(àŒ;Ô˜xÒ'Ð'Ý03µJÔ4`Ñ0aÔ0aÐ-Ø-×4Ò4°jÀ)Ð5LÑMÔMÐMÝ8=½fÐEfÑ>gÔ>gÑ8hÔ8hˆDÐ4Ñ5à TÐ)Ð)Ý',¤~Ô'GˆDÐ#Ñ$Ø�ŠÐuÐuÐuÐu­Ô)H×)MÒ)MÑ)OÔ)OÐuÑuÔuÑvÔvÐvÝÐ!Ð!˜DÐ!Ð!Ð!r   c                 ó*   — |j         |j        |j        fS r)   )r/   r3   Úshared)rc   Úmetadatas     r   Úpack_metadatazHIPBackend.pack_metadata”   s   € àÔØÔØŒOð
ð 	
r   c                 ó.   — 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Ð7Ð7Ð7Ð7Ð7à1°9Ð=Ð=r   c                 óˆ   — t          j        |¦  «         t          j        r!t          j                             |¦  «         d S d S r)   )r	   Úload_dialectsr~   Úinstrumentation)rc   Úctxs     r   r¯   zHIPBackend.load_dialects£   sE   € ÝÔ˜#ÑÔÐÝÔ%ð 	:ÝÔ&×4Ò4°SÑ9Ô9Ð9Ð9Ð9ð	:ð 	:r   c                 ó   — dd l }d}t          | d¦  «        r|                      ¦   «         |k    S t          | |j        ¦  «        r:t          | d¦  «        r*|                      ¦   «                              ¦   «         |k    S dS )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˜Ñ$Ô$ð 	1Ø—=’=‘?”? jÒ0Ð0Ý�c˜5œ<Ñ(Ô(ð 	>­W°SÐ:KÑ-LÔ-Lð 	>Ø×&Ò&Ñ(Ô(×-Ò-Ñ/Ô/°:Ò=Ð=Øˆur   c                 óF   — 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³   s3   € åÔ$ TÑ*Ô*ˆØ�$ˆ;ˆ;ØÐ'¨Ð,Ð-Ñ-ˆCØˆ
r   c                 óŠ   — t          j        | fi |¤Ž}t          j        j        rt
                               | ¦  «        r|dz  }|S )Nr½   )r   Úget_tensor_specializationr
   r	   Úuse_buffer_opsr~   r»   )r¹   ÚkwargsrÀ   s      r   rÂ   z$HIPBackend.get_tensor_specializationº   sJ   € åÔ3°CÐBÐB¸6ÐBÐBˆÝŒ9Ô#ð 	­
×(@Ò(@ÀÑ(EÔ(Eð 	Ø�3‰JˆCØˆ
r   c                 ó*  — t          j        | j        ¦  «        }|                     ¦   «          t          j                             |¦  «         t          j                             |¦  «         t          j
        |j        ¦  «        st          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          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Ô2Ð2ÝÔ ¤Ñ-Ô-ð 	EÝŒK×@Ò@ÀÑDÔDÐDÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò Ñ#Ô#Ð#ÝŒ×)Ò)¨"Ñ-Ô-Ð-ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒ×#Ò# BÑ'Ô'Ð'ÝŒ×$Ò$ RÑ(Ô(Ð(ÝŒ×#Ò# BÑ'Ô'Ð'Ø
�Šˆs�KÑ Ô Ð Øˆ
r   c                 ót  — t          j        | j        ¦  «        }|                     ¦   «          t          j                             |d|j        › �|j        |j	        |j
        ¦  «         |                     | d¦  «         t          j        | j        ¦  «        }|                     ¦   «          d}t          j                             |¦  «         t          j                             ||¦  «         t          j                             |¦  «         t          j                             |¦  «         t"          j        j                             ||j        |j        |j        ¦  «         t          j                             |¦  «         t"          j        j                             |¦  «         t"          j        j                             ||j        ¦  «         t"          j        j                             |¦  «         t"          j        j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t;          |j        ¦  «        }t=          |j        |¦  «        }t"          j        j                             ||j         ¦  «         t"          j        j         !                    |||¦  «         |r*t"          j        j         "                    ||j        ¦  «         t"          j        j         #                    |¦  «         t          j                             |¦  «         |j$         %                    ¦   «         dk    rB|j$         &                    d¦  «        D ]'}t"          j        j         '                    ||¦  «         Œ(t          j                             |¦  «         t          j         (                    |¦  «         tS          |j        ¦  «        rCt"          j        j         *                    |¦  «         t          j                             |¦  «         t"          j        j         +                    |¦  «         |r5|j         dk    r*t"          j        j         ,                    ||j         ¦  «         tZ          j        j.        r�t"          j        j         /                    |¦  «         t          j                             |¦  «         t"          j        j         0                    ||j        tZ          j        j1        tZ          j        j2        ¦  «         t"          j        j         3                    |¦  «         t"          j        j         4                    |¦  «         t          j                             |¦  «         t          j         5                    |¦  «         t          j         6                    |¦  «         |                     | d¦  «         |  7                    ¦   «         |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Ð/D°g´lÐ/DÐ/DÀgÔFWÐY`ÔYjØ+2Ô+;ñ	=ô 	=ð 	=à
�ŠˆsÐ&Ñ'Ô'Ð'ÝŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐØˆÝŒ×#Ò# BÑ'Ô'Ð'ÝŒ×%Ò% b¨'Ñ2Ô2Ð2ÝŒ×4Ò4°RÑ8Ô8Ð8ÝŒ×3Ò3°BÑ7Ô7Ð7ÝŒ
Ô×0Ò0°°W´\À7ÔC_ÐahÔanÑoÔoÐoÝŒ×4Ò4°RÑ8Ô8Ð8ÝŒ
Ô×0Ò0°Ñ4Ô4Ð4ÝŒ
Ô×4Ò4°R¸¼ÑFÔFÐFÝŒ
Ô×7Ò7¸Ñ;Ô;Ð;ÝŒ
Ô×6Ò6°rÑ:Ô:Ð:åŒ×,Ò,¨RÑ0Ô0Ð0ÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×#Ò# BÑ'Ô'Ð'ÝŒ×'Ò'¨Ñ+Ô+Ð+å.¨w¬|Ñ<Ô<ˆÝ9¸'¼,ÈÑWÔWÐåŒ
Ô×-Ò-¨b°'Ô2DÑEÔEÐEÝŒ
Ô×'Ò'¨¨NÐ<NÑOÔOÐOØð 	IÝŒJÔ×6Ò6°r¸7¼<ÑHÔHÐHÝŒ
Ô×4Ò4°RÑ8Ô8Ð8ÝŒ×'Ò'¨Ñ+Ô+Ð+ØÔ ×&Ò&Ñ(Ô(¨FÒ2Ð2ØÔ-×3Ò3°CÑ8Ô8ð Lð L�Ý”
Ô"×AÒAÀ"ÀdÑKÔKÐKÐKÝŒ×4Ò4°RÑ8Ô8Ð8ÝŒ×2Ò2°2Ñ6Ô6Ð6Ý)¨'¬,Ñ7Ô7ð 	=ÝŒJÔ×6Ò6°rÑ:Ô:Ð:ÝŒN×8Ò8¸Ñ<Ô<Ð<ÝŒ
Ô×5Ò5°bÑ9Ô9Ð9Øð 	J 'Ô"4°qÒ"8Ð"8ÝŒJÔ×1Ò1°"°gÔ6HÑIÔIÐIåŒ9Ô#ð 	ÝŒJÔ×8Ò8¸Ñ<Ô<Ð<ÝŒM×+Ò+¨BÑ/Ô/Ð/ÝŒJÔ×8Ò8ØØ”Ý”	Ô,Ý”	Ô?ñ	ô ð õ 	Œ
Ô×-Ò-¨bÑ1Ô1Ð1ÝŒ
Ô×3Ò3°BÑ7Ô7Ð7ÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒ×$Ò$ RÑ(Ô(Ð(Ø
�Šˆs�LÑ!Ô!Ð!Ø&)×&AÒ&AÑ&CÔ&CˆÐ"Ñ#Øˆ
r   c                 ó¸  — | }t          j        |j        ¦  «        }|                     ¦   «          t          j                             |¦  «         t          j                             |¦  «         t          j         	                    |¦  «         t          j
                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j        j                             |¦  «         t          j                             |¦  «         |                     |d¦  «         |                     ¦   «         |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Ô3Ð3ÝŒ×Ò˜rÑ"Ô"Ð"ÝŒ×&Ò& rÑ*Ô*Ð*ÝŒ×&Ò& rÑ*Ô*Ð*ÝŒ×7Ò7¸Ñ;Ô;Ð;ÝŒ
Ô×,Ò,¨RÑ0Ô0Ð0ÝŒ×/Ò/°Ñ3Ô3Ð3à
�ŠˆsÐ$Ñ%Ô%Ð%Ø&)×&AÒ&AÑ&CÔ&CˆÐ"Ñ#Øˆ
r   c                 ó¢  ‡— | }t          j        |j        ¦  «        }|                     ¦   «          t          j        j                             ||j        ¦  «         t          j        j         	                    ||j        ¦  «         t
          j
                             |¦  «         t
          j                             |¦  «         t
          j
                             |¦  «         t          j        j                             |¦  «         t           j        r&t           j                             d||j        ¦  «         d}t          j        j                             ||j        |¦  «         t          j        j                             ||j        ¦  «         t
          j                             |¦  «         t
          j                             |¦  «         t
          j
                             |¦  «         t
          j
                             |¦  «         t
          j                             |¦  «         t
          j                             |¦  «         t
          j                             |¦  «         |j                             ¦   «         dk    r0t          j        j                             ||j        |j        ¦  «         t           j        r&t           j                             d||j        ¦  «         t>          j         j!        s0t>          j         j"        st
          j#         $                    |¦  «         t          j        j         %                    ||j        |¦  «         | &                    |d¦  «         t>          j         j"        rÕt>          j         j!        sbt          j        |j        ¦  «        }|                     ¦   «          t
          j#         $                    |¦  «         | &                    |d¦  «         t          j        |j        ¦  «        }|                     ¦   «          t
          j#         '                    |¦  «         | &                    |d¦  «         tQ          j)        ¦   «          tQ          j        ¦   «         }tQ          j*        ||¦  «        Št	          j+        ‰¦  «         d}t>          j         j,        rd	}tQ          j-        ‰t          j.        |j        |¦  «         t	          j/        ‰|j        ¦  «         t	          j0        ‰d
¦  «         t	          j1        ‰dd¦  «         t	          j1        ‰dd¦  «         t	          j1        ‰dd¦  «         t	          j1        ‰d|j2        dk    ¦  «         d„ ‰ 3                    ¦   «         D ¦   «         }|d          4                    t          j5        ¦  «         |d         }	|d          6                    d|	› d�¦  «         |j7        }
|  8                    d¦  «        }|�|}
|d          6                    dd|
|j2        z  › �¦  «         d|j         9                    d¦  «        v r|d          6                    dd¦  «         |d          6                    dd¦  «         |d          6                    d|j:        › d |j:        › �¦  «         |j;        rd!nd"}|d          6                    d#|¦  «         t>          j         j,        r5|d          <                    d	¦  «         |d          =                    ¦   «          |j        d$k    r<t>          j        j>        s+t>          j        j?        st	          j@        |d         ¦  «         t>          j         j,        rgtƒ          t„          ¦  «        jC        d%z  }t‰          |d&z  ¦  «        t‰          |d'z  ¦  «        t‰          |d(z  ¦  «        g}tQ          jE        ‰|¦  «         nB|jF        r;ˆfd)„|jF        D ¦   «         }t�          |¦  «        dk    rtQ          jE        ‰|¦  «         tQ          jH        ‰tP          jI        |j        dg |jJ        ¦  «         t	          jK        |j        ¦  «        rQ|d          L                    d*¦  «         |d          L                    d+¦  «         |d          L                    d,¦  «         t>          j        jM        rt	          jN        |d         ¦  «         |
|d-<   |  8                    d.¦  «        |d/<   |  8                    d0¦  «        pd|d1<   |  8                    d2¦  «        pd3|d4<   t	          jO        ‰¦  «         t	          jP        ‰¦  «         t‰          ‰¦  «        S )5NÚ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   c                 ó:   — g | ]}|                      ¦   «         °|‘ŒS r   )Úis_declaration)rk   Úfns     r   rn   z(HIPBackend.make_llir.<locals>.<listcomp>~  s)   € ÐPÐPÐP�b¸B×<MÒ<MÑ<OÔ<OÐPˆrÐPÐPÐPr   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.bcc                 óD   •— g | ]\  }}t          j        ‰|¦  «        ¯|‘ŒS r   )r	   Úneed_extern_lib)rk   rl   ÚpathÚllvm_mods      €r   rn   z(HIPBackend.make_llir.<locals>.<listcomp>­  s1   ø€ ÐiÐiÐi™l˜t TÅSÔEXÐYaÐcgÑEhÔEhÐi�TÐiÐiÐir   z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)Qr   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Ú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   Ú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È   Útarget_featuresÚfnsÚcluster_dimÚtotal_warps_numÚtotal_num_warpsÚdenormal_modere   Úpathsr  s                  @r   r  zHIPBackend.make_llir(  s  ø€ àˆåŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐÝŒ
Ô×6Ò6°r¸7¼<ÑHÔHÐHÝŒ
Ô×7Ò7¸¸G¼LÑIÔIÐIÝŒ×$Ò$ RÑ(Ô(Ð(ÝŒ× Ò  Ñ$Ô$Ð$ÝŒ×*Ò*¨2Ñ.Ô.Ð.åŒ
Ô×5Ò5°bÑ9Ô9Ð9åÔ%ð 	SÝÔ&×,Ò,Ð-@À"ÀcÄkÑRÔRÐRð ˆ	ÝŒ
Ô×(Ò(¨¨W¬\¸9ÑEÔEÐEÝŒ
Ô×6Ò6°r¸7¼<ÑHÔHÐHÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò˜bÑ!Ô!Ð!åŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×*Ò*¨2Ñ.Ô.Ð.ÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒ×$Ò$ RÑ(Ô(Ð(àÔ ×&Ò&Ñ(Ô(¨FÒ2Ð2ÝŒJÔ×<Ò<¸RÀÄÈwÔOaÑbÔbÐbõ Ô%ð 	PÝÔ&×,Ò,Ð-=¸rÀ3Ä;ÑOÔOÐOåÔ Ô2ð 	+½5Ô;LÔ;oð 	+ÝŒM×&Ò& rÑ*Ô*Ð*åŒ
Ô×5Ò5°b¸'¼,È	ÑRÔRÐRØ
�Šˆs�KÑ Ô Ð åÔÔ?ð 	HåÔ$Ô6ð ;Ý”_ S¤[Ñ1Ô1�Ø—’Ñ!Ô!Ð!Ý”×*Ò*¨2Ñ.Ô.Ð.Ø—’�sÐ9Ñ:Ô:Ð:õ ” ¤Ñ-Ô-ˆBØ�OŠOÑÔÐÝŒM×/Ò/°Ñ3Ô3Ð3Ø�FŠF�3ÐFÑGÔGÐGõ 	ÔÑÔÐÝ”,‘.”.ˆÝ”> # wÑ/Ô/ˆÝÔ  Ñ*Ô*Ð*ØˆÝÔÔ(ð 	'Ø&ˆOÝÔ˜x­Ô):¸G¼LÈ/ÑZÔZÐZõ 	Ô˜H g¤lÑ3Ô3Ð3ÝÔ˜H cÑ*Ô*Ð*ÝÔ% hÐ0HÈ%ÑPÔPÐPÝÔ% hÐ0QÐSWÑXÔXÐXÝÔ% hÐ0HÈ%ÑPÔPÐPÝÔ% hÐ0HÈ'ÔJ[Ð_aÒJaÑbÔbÐbð QÐP˜H×2Ò2Ñ4Ô4ÐPÑPÔPˆàˆAŒ×Ò¥Ô >Ñ?Ô?Ð?Ø˜zÔ*ˆØˆAŒ×ÒÐ0°[Ð2FÐ2FÐ2FÑGÔGÐGà!Ô+ˆØ×*Ò*Ð+@ÑAÔAˆØÐ&Ø-ˆOØˆAŒ×ÒÐ8Ð:b¸ÈwÔO`Ñ?`Ð:bÐ:bÑcÔcÐcØ# wÔ'<×'BÒ'BÀ3Ñ'GÔ'GÐGÐGØ�ŒF×ÒÐ6¸ÑHÔHÐHØˆAŒ×ÒÐ4°fÑ=Ô=Ð=ð 	ˆAŒ×ÒÐ0°WÔ5IÐ2cÐ2cÈWÔMaÐ2cÐ2cÑdÔdÐdØ+2Ô+EÐQ˜˜È6ˆØˆAŒ×ÒÐ1°=ÑAÔAÐAÝÔÔ(ð 	&Ø�ŒF×(Ò(¨Ñ2Ô2Ð2Ø�ŒF×#Ò#Ñ%Ô%Ð%ð Œ<˜9Ò$Ð$­e¬iÔ.@Ð$ÅEÄIÔDVÐ$ÝÔ$ S¨¤VÑ,Ô,Ð,åÔÔ(ð 	7Ý!¥(™^œ^Ô2°UÑ:ˆNå�N \Ñ1Ñ2Ô2Ý�N YÑ.Ñ/Ô/Ý�N YÑ.Ñ/Ô/ðˆEõ
 Ô! (¨EÑ2Ô2Ð2Ð2ØÔ ð 	7ØiÐiÐiÐi¨gÔ.AÐiÑiÔiˆEÝ�5‰zŒz˜AŠ~ˆ~ÝÔ% h°Ñ6Ô6Ð6åÔ˜X¥tÔ'7¸¼ÀrÈ2ÈwÔOgÑhÔhÐhõ Ô$ W¤\Ñ2Ô2ð 	>Ø�ŒF×!Ò!Ð"<Ñ=Ô=Ð=Ø�ŒF×!Ò!Ð"<Ñ=Ô=Ð=Ø�ŒF×!Ò!Ð"<Ñ=Ô=Ð=åŒ9Ô*ð 	<ÝÔ3°C¸´FÑ;Ô;Ð;ð !0ˆ�ÑØ ×-Ò-¨lÑ;Ô;ˆ�ÑØ+.×+;Ò+;Ð<]Ñ+^Ô+^Ð+cÐbcˆÐ'Ñ(Ø,/×,<Ò,<Ð=cÑ,dÔ,dÐ,iÐhiˆÐ(Ñ)åÔ$ XÑ.Ô.Ð.õ 	Ô  Ñ*Ô*Ð*Ý�8‰}Œ}Ðr   c           	      ón  — t          j        d| ¦  «        }t          |¦  «        dk    sJ ‚|d         |d<   g }d|j        v rdnd}t	          j        |                      d¦  «        ¦  «                             ¦   «         }|d         d	z   |z   }t          j	        | t          j        |j        |||j        |¦  «        }t          j        | t          j        |j        |||j        |¦  «         t          j
        j        r_t          j        t"          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   rl   Úgfx11z-real-true16rJ   ro   rh   z.txtFz!// -----// AMDGCN Dump //----- //)ÚreÚfindallrD  r!   rr   rs   rt   ru   r   Útranslate_to_mirr	   r5  rB   Údump_sched_dagr
   r@  Útranslate_mir_to_asmÚosr  rp   Ú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ÔWˆÝ�5‰zŒz˜QŠˆˆˆØ  œ8ˆ�ÑàˆØ%,°´Ð%<Ð%<�>�>À"ˆÝ”. §¢¨GÑ!4Ô!4Ñ5Ô5×?Ò?ÑAÔAˆØ˜Q”x #‘~¨Ñ/ˆÝÔ! #¥sÔ'8¸'¼,ÈÐRWÐY`ÔYqØ".ñ0ô 0ˆåÔ˜C¥Ô!2°G´LÀ(ÈEÐSZÔSkØ(ñ	*ô 	*ð 	*åŒ9Ôð 	LÝÔ.­r¬w¯|ª|½E¼IÔ<NØ<HÈ6Ñ<Qñ0Sô 0SÝTWÔTeÐgnÔgsØ/7¸ÀÔ@XÐZ_ñaô aˆFˆFõ Ô*¨3µÔ0AÀ7Ä<ÐQYÐ[`Ø+2Ô+CÀUñLô LˆFåŒ9Ô ð 	ÝÐ5Ñ6Ô6Ð6Ý�&‰MŒMˆMØˆr   c                 ór  — 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 }| 
                    |¦  «         d d d ¦  «         n# 1 swxY w Y   t	          j        |j	        |j	        ¦  «         d d d ¦  «         n# 1 swxY w Y   t          |j	        d¦  «        5 }|                     ¦   «         }	d d d ¦  «         n# 1 swxY w Y   d d d ¦  «         n# 1 swxY w Y   |	S )NrJ   r  rV  z,-real-true16ÚwbÚrb)r
   r)  r3  r!   r	   Úassemble_amdgcnÚtempfileÚNamedTemporaryFileÚopenrl   ÚwriteÚ
link_hsacoÚread)
r  r¥   r�   rN  r„   Útmp_outÚtmp_inÚfd_inÚfd_outrÀ   s
             r   Ú
make_hsacozHIPBackend.make_hsacoè  s  € àˆÝÔÔ(ð 	'Ø&ˆOØ�g”lÐ"Ð"Ø˜Ñ.ˆOÝÔ# C¨¬°ÑGÔGˆÝÔ(Ñ*Ô*ð 	$¨gÝÔ,Ñ.Ô.ð :°&Ý˜&œ+ tÑ,Ô,ð '°Ø—K’K Ñ&Ô&Ð&ð'ð 'ð 'ñ 'ô 'ð 'ð 'ð 'ð 'ð 'ð 'øøøð 'ð 'ð 'ð 'å”˜vœ{¨G¬LÑ9Ô9Ð9ð:ð :ð :ñ :ô :ð :ð :ð :ð :ð :ð :øøøð :ð :ð :ð :õ �g”l DÑ)Ô)ð $¨VØ—k’k‘m”m�ð$ð $ð $ñ $ô $ð $ð $ð $ð $ð $ð $øøøð $ð $ð $ð $ð	$ð 	$ð 	$ñ 	$ô 	$ð 	$ð 	$ð 	$ð 	$ð 	$ð 	$øøøð 	$ð 	$ð 	$ð 	$ð ˆ
s~   ÁD,Á&CÁ<BÂCÂB"Â"CÂ%B"Â&"CÃD,ÃC	ÃD,ÃC	ÃD,Ã4DÄ	D,ÄD	ÄD,ÄD	ÄD,Ä,D0Ä3D0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                             ‰ |‰|d ¦  «         d S d S )Nc                 ó2   •— ‰                      | |‰¦  «        S r)   )rÆ   ©r  r¥   r�   rc   s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>û  s   ø€ °4·>²>À#ÀxÐQXÑ3YÔ3Y€ r   rÌ   c                 ó2   •— ‰                      | |‰¦  «        S r)   )rÝ   rx  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>ü  s   ø€ °D·O²OÀCÈÐSZÑ4[Ô4[€ r   Úttgirc                 ó2   •— ‰                      | |‰¦  «        S r)   )r   rx  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>þ  s   ø€ °D×4GÒ4GÈÈXÐW^Ñ4_Ô4_€ r   c                 ó2   •— ‰                      | |‰¦  «        S r)   )r  rx  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>ÿ  s   ø€ ¨t¯~ª~¸cÀ8ÈWÑ/UÔ/U€ r   Úllirc                 ó2   •— ‰                      | |‰¦  «        S r)   )rf  rx  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>   s   ø€ °×1AÒ1AÀ#ÀxÐQXÑ1YÔ1Y€ r   re  c                 ó2   •— ‰                      | |‰¦  «        S r)   )ru  rx  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>  s   ø€ °·²ÀÀXÈwÑ0WÔ0W€ r   r„   )r   ÚTRITONÚGLUONr
   r•   Úadd_stages_inspection_hook)rc   Ústagesr�   r�   s   ` ` r   Ú
add_stageszHIPBackend.add_stagesù  sØ   øø€ Ø•x”Ò&Ð&ØYÐYÐYÐYÐYˆF�6‰NØ[Ð[Ð[Ð[Ð[ˆF�7‰OˆOØ�œÒ'Ð'Ø_Ð_Ð_Ð_Ð_ˆF�7‰OØUÐUÐUÐUÐUˆˆv‰ØYÐYÐYÐYÐYˆˆxÑØWÐWÐWÐWÐWˆˆw‰ÝŒ=Ô3Ð?ÝŒM×4Ò4°T¸6À7ÈHÐVZÑ[Ô[Ð[Ð[Ð[ð @Ð?r   c                 ó   — | j         › S r)   r   )rc   s    r   rw   zHIPBackend.hash  s   € à”+ÐÐr   )!rx   ry   rz   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  rf  ru  r„  Ú	functoolsÚ	lru_cacherw   Ú__classcell__)r‰   s   @r   r~   r~   k   s3  ø€ € € € € Ø€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   r\  rr   rk  rW  rˆ  r[   Úpathlibr   r   r#   r'   r*   r-   r~   r   r   r   ú<module>r’     s½  ðØ EÐ EÐ EÐ EÐ EÐ EÐ EÐ EÐ EÐ EØ 5Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5Ø Ð Ð Ð Ð Ð Ø !Ð !Ð !Ð !Ð !Ð !Ø #Ð #Ð #Ð #Ð #Ð #Ð #Ð #Ð #Ð #Ø Ð Ð Ð Ð Ð Ø 	€	€	€	Ø €€€Ø €€€Ø 	€	€	€	Ø Ð Ð Ð Ø €€€Ø Ð Ð Ð Ð Ð ð0˜Yð 0ð 0ð 0ð 0ðRð Rð Rð
rð rð rð$ð $ð $ð €�$ÐÑÔðD?ð D?ð D?ð D?ð D?ñ D?ô D?ñ ÔðD?ðN\ ð \ ð \ ð \ ð \ �ñ \ ô \ ð \ ð \ ð \ r   