§
    8˜ºhþM  ã                   ó  — d dl mZmZmZ d dlmZmZmZmZ d dl	m
Z
 d dlmZ d dlmZmZmZ d dlmZ d dlZd dlZd dlZd dlZd dlZd dlmZ d	efd
„Zd„ Zd„ Z ed¬¦  «         G d„ d¦  «        ¦   «         Z G d„ de¦  «        Z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     úd/var/www/html/services/rag-service/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                 óR   — t           j        j        €| dk    nt           j        j        S ©NÚgfx942)r
   r	   Úuse_block_pingpong©Úarchs    r   Úis_pingpong_schedule_enabledr$      s!   € Ý!&¤Ô!=Ð!EˆD�HÒÐÍ5Ì9ÔKgÐgr   c                 óR   — t           j        j        €| dk    nt           j        j        S r   )r
   r	   Úuse_in_thread_transposer"   s    r   Úis_in_thread_transpose_enabledr'      s#   € Ý!&¤Ô!BÐ!JˆD�HÒÐÕPUÔPYÔPqÐqr   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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 dS )&Ú
HIPOptionsé   Ú	num_warpsr   Úwaves_per_eué   Ú
num_stagesÚnum_ctasNÚextern_libsr   Úcluster_dimsFÚdebugTÚsanitize_overflowr#   )Úfp8e5Úsupported_fp8_dtypesr   Ú!deprecated_fp8_dot_operand_dtypesÚieeeÚdefault_dot_input_precision)r8   Úallowed_dot_input_precisionsÚenable_fp_fusionÚlaunch_cooperative_gridr   Úmatrix_instr_nonkdimÚkpackÚallow_flush_denormÚmax_num_imprecise_acc_defaultÚhipÚbackend_nameÚnoneÚschedule_hintc                 óB  — 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    r| j        dk    s
J 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 2Úgfx950zgfx950 only accepts kpack == 1Úlib)ÚocmlÚocklz.bcr1   )Úintr#   ÚobjectÚ__setattr__r,   r>   r   Ú__file__Úparentr1   ÚdictÚstrÚtupleÚitems)ÚselfÚ	gfx_majorrK   Údefault_libdirr1   rM   s         r   Ú__post_init__zHIPOptions.__post_init__E   s1  € Ý˜œ	 ! B $œÑ(Ô(ˆ	Ø# rš/˜/�B�B¨rˆ	Ý×Ò˜4 ¨iÑ8Ô8Ð8ØŒ~ Ò!Ð! t¤~¸¼È!Ñ9KÑ'LÐQRÒ&RÐ&RÐ&RØ/ñ 'SÔ&RÐ&Rð Œ9˜Ò Ð Ø”: ’?�?�?Ð$D‘?”?�?å�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>V   s&   € ÐOÐOÐO©I¨D°#˜4˜˜ #˜˜ÐOÐOÐOr   úutf-8)ÚjoinÚ__dict__rX   ÚhashlibÚsha256ÚencodeÚ	hexdigest)rY   Úkeys     r   ÚhashzHIPOptions.hashU   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,   rP   Ú__annotations__r-   r/   r0   r1   rU   r2   rW   r3   Úboolr4   r#   rV   r6   r   r7   r9   r:   r;   r<   r=   r>   r?   r@   rB   rD   r\   rm   r   r   r   r*   r*      sœ  € € € € € € à€IˆsÐÐÑØ€L�#ÐÐÑØ€J�ÐÐÑØ€HˆcÐÐÑØ€K�ÐÐÑØ#€L�%Ð#Ð#Ñ#Ø€Eˆ4ÐÐÑØ"Ð�tÐ"Ð"Ñ"Ø€Dˆ#ÐÐÑØ'2Ð˜% œ*Ð2Ð2Ñ2Ø46Ð% u¨S¤zÐ6Ð6Ñ6Ø'-Ð Ð-Ð-Ñ-Ø/9Ð  %¨¤*Ð9Ð9Ñ9Ø!Ð�dÐ!Ð!Ñ!Ø$)Ð˜TÐ)Ð)Ñ)Ø !Ð˜#Ð!Ð!Ñ!Ø€Eˆ3€N€N�NØ$Ð˜Ð$Ð$Ñ$Ø)*Ð! 3Ð*Ð*Ñ*Ø€L�#ÐÐÑð"  €M�3ÐÐÑðLð Lð Lð ?ð ?ð ?ð ?ð ?r   r*   c                   óœ  ‡ — e Zd Zedefd„¦   «         Zdeddfˆ fd„Zdefd„Zde	fd„Z
d„ Zd	„ Zdeeef         fd
„Zd„ Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Zd„ Z ej        ¦   «         d„ ¦   «         Zˆ xZS )Ú
HIPBackendr   c                 ó   — | j         dk    S )NrA   )Úbackendr   s    r   Úsupports_targetzHIPBackend.supports_target\   s   € àŒ~ Ò&Ð&r   ÚreturnNc                 ó�   •— t          ¦   «                              |¦  «         t          |j        t          ¦  «        sJ ‚d| _        d S )NÚhsaco)ÚsuperÚ__init__Ú
isinstancer#   rV   Ú
binary_ext)rY   r   Ú	__class__s     €r   r|   zHIPBackend.__init__`   s>   ø€ Ý‰Œ×Ò˜Ñ Ô Ð Ý˜&œ+¥sÑ+Ô+Ð+Ð+Ð+Ø!ˆŒˆˆr   c                 ó   — d|j         › �S )Núhip:r"   ©rY   Úoptionss     r   Úget_target_namezHIPBackend.get_target_namee   s   € Ø$�g”lÐ$Ð$Ð$r   c                 ó"  ‡— dt           j        j        p| j        j        i}| j        j        dk    rNt          t          j        ¦  «        }|                     dh¦  «         t          t          |¦  «        ¦  «        |d<   d‰vr­t          t          j        ¦  «        }| j        j        dk    r|                     h d£¦  «         nM| j        j        dk    r|                     dd	h¦  «         n%d
| j        j        v r|                     dd	h¦  «         t          t          |¦  «        ¦  «        |d<   d‰vrt           j        j        |d<   |                     ˆfd„t          j                             ¦   «         D ¦   «         ¦  «         t          di |¤ŽS )Nr#   r    Útf32r:   r6   >   Úfp8e4b8Úfp8e4nvÚfp8e5b16rL   rˆ   r5   Úgfx12r;   c                 ó:   •— i | ]}|‰v ¯‰|         ®|‰|         “ŒS ©Nr   )ra   ÚkÚoptss     €r   ú
<dictcomp>z,HIPBackend.parse_options.<locals>.<dictcomp>}   s:   ø€ ð ;ð ;ð ; AØ˜T˜	˜	 d¨1¤gÐ&9ð ˜˜QœØ&9Ð&9Ð&9r   r   )r
   ÚruntimeÚoverride_archr   r#   Úsetr*   r:   ÚupdaterW   Úsortedr6   ÚlanguageÚdefault_fp_fusionÚ__dataclass_fields__Úkeys)rY   rŽ   Úargsr:   r6   s    `   r   Úparse_optionszHIPBackend.parse_optionsh   s¥  ø€ Ø�œÔ3ÐG°t´{Ô7GÐHˆð Œ;Ô˜xÒ'Ð'Ý+.­zÔ/VÑ+WÔ+WÐ(Ø(×/Ò/°°Ñ9Ô9Ð9Ý38½Ð@\Ñ9]Ô9]Ñ3^Ô3^ˆDÐ/Ñ0à!¨Ð-Ð-Ý#&¥zÔ'FÑ#GÔ#GÐ ØŒ{Ô 8Ò+Ð+Ø$×+Ò+Ð,NÐ,NÐ,NÑOÔOÐOÐOØ”Ô! XÒ-Ð-Ø$×+Ò+¨Y¸Ð,@ÑAÔAÐAÐAØ˜DœKÔ,Ð,Ð,Ø$×+Ò+¨Y¸Ð,@ÑAÔAÐAÝ+0µÐ8LÑ1MÔ1MÑ+NÔ+NˆDÐ'Ñ(à TÐ)Ð)Ý',¤~Ô'GˆDÐ#Ñ$Ø�Šð ;ð ;ð ;ð ;­Ô)H×)MÒ)MÑ)OÔ)Oð ;ñ ;ô ;ñ 	<ô 	<ð 	<åÐ!Ð!˜DÐ!Ð!Ð!r   c                 ór   — |j         |j        |j        |j        d         |j        d         |j        d         fS )Nr   r   r.   )r,   r0   Úsharedr2   )rY   Úmetadatas     r   Úpack_metadatazHIPBackend.pack_metadata�   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£   )rY   r£   s     r   Úget_module_mapzHIPBackend.get_module_mapŽ   s   € Ø7Ð7Ð7Ð7Ð7Ð7à1°9Ð=Ð=r   c                 ó.   — t          j        |¦  «         d S rŒ   )r	   Úload_dialects)rY   Úctxs     r   r§   zHIPBackend.load_dialects“   s   € ÝÔ˜#ÑÔÐÐÐ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_rangerI   )r   Ú
parse_attr)ÚdescÚrets     r   rµ   zHIPBackend.parse_attr¡   s3   € åÔ$ TÑ*Ô*ˆØ�$ˆ;ˆ;ØÐ'¨Ð,Ð-Ñ-ˆCØˆ
r   c                 ó˜   — t          j        | |fi |¤Ž}t          j        j        r%|dk    rt
                               | ¦  «        r|dz  }|S )NÚtensorr´   )r   Úget_arg_specializationr
   r	   Úuse_buffer_opsrt   r²   )r°   ÚtyÚkwargsr·   s       r   rº   z!HIPBackend.get_arg_specialization¨   sT   € åÔ0°°bÐCÐC¸FÐCÐCˆõ Œ9Ô#ð 	¨¨hª¨½:×;SÒ;SÐTWÑ;XÔ;X¨Ø�3‰JˆCØˆ
r   c                  óˆ  — t           j        j        } | �%t          | ¦  «        }|                     ¦   «         r|S t          t
          ¦  «        j        dz  }|                     ¦   «         r|S t          d¦  «        }|                     ¦   «         r|S t          d¦  «        }|                     ¦   «         r|S t          d¦  «        ‚)Nzllvm/bin/ld.lldz/opt/rocm/llvm/bin/ld.lldz/usr/bin/ld.lldzWROCm linker /opt/rocm/llvm/bin/ld.lld not found. Set 'TRITON_HIP_LLD_PATH' to its path.)r
   r	   Úlld_pathr   Úis_filerS   rT   Ú	Exception)Úlld_env_pathÚllds     r   Úpath_to_rocm_lldzHIPBackend.path_to_rocm_lld±   s´   € õ ”yÔ)ˆØÐ#Ý�|Ñ$Ô$ˆCØ�{Š{‰}Œ}ð Ø�
å•8‰nŒnÔ#Ð&7Ñ7ˆØ�;Š;‰=Œ=ð 	ØˆJÝÐ.Ñ/Ô/ˆØ�;Š;‰=Œ=ð 	ØˆJÝÐ$Ñ%Ô%ˆØ�;Š;‰=Œ=ð 	ØˆJÝÐqÑrÔrÐrr   c                 óö  — t          j        | j        ¦  «        }|                     ¦   «          t          j                             |¦  «         t          j                             |¦  «         t          j         	                    |¦  «         t          j         
                    |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         |                     | ¦  «         | S rŒ   )r   Úpass_managerÚcontextÚenable_debugr   ÚcommonÚadd_inlinerÚttirÚadd_rewrite_tensor_pointerÚ(add_rewrite_tensor_descriptor_to_pointerÚadd_canonicalizerÚadd_combineÚadd_reorder_broadcastÚadd_cseÚadd_triton_licmÚadd_symbol_dceÚadd_loop_unrollÚrun)Úmodr�   rƒ   Úpms       r   Ú	make_ttirzHIPBackend.make_ttirÅ   s  € åŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐÝŒ×!Ò! "Ñ%Ô%Ð%ÝŒ×.Ò.¨rÑ2Ô2Ð2ÝŒ×<Ò<¸RÑ@Ô@Ð@ÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò Ñ#Ô#Ð#ÝŒ×)Ò)¨"Ñ-Ô-Ð-ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒ×#Ò# BÑ'Ô'Ð'ÝŒ×$Ò$ RÑ(Ô(Ð(ÝŒ×#Ò# BÑ'Ô'Ð'Ø
�Šˆs‰ŒˆØˆ
r   c                 óN  — t          j        | j        ¦  «        }|                     ¦   «          t          j                             |d|j        › �|j        |j	        |j
        ¦  «         |                     | ¦  «         t          j        | j        ¦  «        }|                     ¦   «          t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t           j        j                             ||j        |j        |j        ¦  «         t          j                             |¦  «         t           j        j                             |¦  «         t          j                             |d¦  «         t           j        j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t          j                             |¦  «         t6          j        j        }t6          j        j        }t6          j        j        }|j        dk    rdx}}t           j        j                              ||j!        |||¦  «         |r*t           j        j         "                    ||j        ¦  «         t          j                             |¦  «         |j         #                    ¦   «         dk    r*t           j        j         $                    ||j        ¦  «         t          j                             |d¦  «         t          j                             |¦  «         t          j         %                    |¦  «         tM          |j        ¦  «        rCt           j        j         '                    |¦  «         t          j                             |¦  «         t           j        j         (                    |¦  «         tS          |j        ¦  «        }|r5|j!        dk    r*t           j        j         *                    ||j!        ¦  «         t6          j        j+        rmt           j        j         ,                    |¦  «         t          j                             |¦  «         t           j        j         -                    ||j        ¦  «         t           j        j         .                    |¦  «         t          j                             |¦  «         t          j         /                    |¦  «         t          j         0                    |¦  «         |r*t           j        j         1                    ||j        ¦  «         |                     | ¦  «         | S )Nr�   Tzlocal-prefetchr   rC   r.   )2r   rÆ   rÇ   rÈ   r   rË   Úadd_convert_to_ttgpuirr#   r,   rK   r0   rÕ   ÚttgpuirÚadd_coalesceÚadd_remove_layout_conversionsÚadd_optimize_thread_localityr	   Úadd_accelerate_matmulr=   r>   Úadd_optimize_epilogueÚadd_optimize_dot_operandsÚadd_hoist_layout_conversionsÚadd_fuse_nested_loopsrÉ   rÎ   rÒ   r
   Úglobal_prefetchÚlocal_prefetchÚuse_async_copyrD   Úadd_stream_pipeliner/   Úadd_coalesce_async_copyÚlowerÚinsert_instruction_sched_hintsÚadd_reduce_data_duplicationr'   Úadd_in_thread_transposeÚadd_reorder_instructionsr$   Úadd_block_pingpongr»   Úadd_canonicalize_pointersÚadd_convert_to_buffer_opsÚadd_fold_true_cmpirÑ   rÓ   Úadd_update_async_wait_count)rÖ   r�   rƒ   r×   rä   rå   ræ   r!   s           r   Ú
make_ttgirzHIPBackend.make_ttgirÖ   sS  € åŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐÝŒ×*Ò*¨2Ð/D°g´lÐ/DÐ/DÀgÔFWÐY`ÔYjØ+2Ô+;ñ	=ô 	=ð 	=à
�Šˆs‰ŒˆÝŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐÝŒ×#Ò# BÑ'Ô'Ð'ÝŒ×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ÝŒ×0Ò0°°TÑ:Ô:Ð:ÝŒ
Ô×7Ò7¸Ñ;Ô;Ð;åŒ×,Ò,¨RÑ0Ô0Ð0ÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×#Ò# BÑ'Ô'Ð'ÝŒ×'Ò'¨Ñ+Ô+Ð+åœ)Ô3ˆÝœÔ1ˆÝœÔ1ˆð Ô Ð$4Ò4Ð4Ø/0Ð0ˆO˜nåŒ
Ô×.Ò.¨r°7Ô3EÀÐXfÐhvÑwÔwÐwØð 	IÝŒJÔ×6Ò6°r¸7¼<ÑHÔHÐHÝŒ×'Ò'¨Ñ+Ô+Ð+ØÔ ×&Ò&Ñ(Ô(¨FÒ2Ð2ÝŒJÔ×=Ò=¸bÀ'ÔBWÑXÔXÐXÝŒ×0Ò0°°TÑ:Ô:Ð:ÝŒ×4Ò4°RÑ8Ô8Ð8ÝŒ×2Ò2°2Ñ6Ô6Ð6Ý)¨'¬,Ñ7Ô7ð 	=ÝŒJÔ×6Ò6°rÑ:Ô:Ð:ÝŒN×8Ò8¸Ñ<Ô<Ð<ÝŒ
Ô×3Ò3°BÑ7Ô7Ð7Ý9¸'¼,ÑGÔGÐØð 	J 'Ô"4¸Ò"9Ð"9ÝŒJÔ×1Ò1°"°gÔ6HÑIÔIÐIåŒ9Ô#ð 	KÝŒJÔ×8Ò8¸Ñ<Ô<Ð<ÝŒM×+Ò+¨BÑ/Ô/Ð/ÝŒJÔ×8Ò8¸¸W¼\ÑJÔJÐJåŒ
Ô×-Ò-¨bÑ1Ô1Ð1ÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒ×$Ò$ RÑ(Ô(Ð(Øð 	MÝŒJÔ×:Ò:¸2¸w¼|ÑLÔLÐLØ
�Šˆs‰ŒˆØˆ
r   c                 óÄ  — | }t          j        |j        ¦  «        }|                     ¦   «          t          j                             |¦  «         t          j                             |¦  «         t          j	         
                    |¦  «         t          j                             |¦  «         t          j                             |¦  «         |                     |¦  «         |S rŒ   )r   rÆ   rÇ   rÈ   r   rÛ   rÊ   rÉ   Úadd_sccprË   Úadd_loop_aware_cserÎ   Ú add_combine_tensor_select_and_ifrÕ   )Úsrcr�   rƒ   rÖ   r×   s        r   Ú	ttgir_optzHIPBackend.ttgir_opt  sª   € àˆÝŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐåŒ×"Ò" 2Ñ&Ô&Ð&ÝŒ×Ò˜rÑ"Ô"Ð"ÝŒ×&Ò& rÑ*Ô*Ð*ÝŒ×(Ò(¨Ñ,Ô,Ð,ÝŒ×7Ò7¸Ñ;Ô;Ð;à
�Šˆs‰ŒˆØˆ
r   c                 óð  ‡— | }t          j        |j        ¦  «        }|                     ¦   «          d}t          j        j                             ||j        |¦  «         t
          j	         
                    |¦  «         t
          j	                             |¦  «         t
          j                             |¦  «         d}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        ¦  «         t0          j        j        st
          j                             |¦  «         t          j        j                             ||¦  «         |                     |¦  «         t?          j         ¦   «          t?          j        ¦   «         }t?          j!        ||¦  «        Št	          j"        ‰¦  «         d}t0          j        j#        rd}t?          j$        ‰t          j%        |j        |¦  «         t	          j&        ‰|j        ¦  «         t	          j'        ‰d¦  «         t	          j(        ‰dd¦  «         t	          j(        ‰d	d¦  «         t	          j(        ‰d
d¦  «         t	          j(        ‰d|j)        dk    ¦  «         d„ ‰ *                    ¦   «         D ¦   «         }	|	d          +                    t          j,        ¦  «         |	d          -                    dd|j.        |j)        z  › �¦  «         |	d          -                    d|j/        › ¦  «         |j0        rdnd}
|	d          -                    d|
¦  «         t0          j        j#        r5|	d          1                    d¦  «         |	d          2                    ¦   «          t	          j3        |	d         ¦  «         t0          j        j#        rgti          tj          ¦  «        j6        dz  }to          |dz  ¦  «        to          |dz  ¦  «        to          |dz  ¦  «        g}t?          j8        ‰|¦  «         n/|j9        r(ˆfd„|j9        D ¦   «         }t?          j8        ‰|¦  «         t?          j:        ‰t>          j;        |j        dg |j<        ¦  «         t0          j        j=        rt	          j>        |	d         ¦  «         |  ?                    d¦  «        |d<   t	          j@        ‰¦  «         t	          jA        ‰¦  «         to          ‰¦  «        S )Nr   TrC   Ú ú+xnackiô  Ú__oclc_finite_only_optFÚ__oclc_correctly_rounded_sqrt32Ú__oclc_unsafe_math_optÚ__oclc_wavefrontsize64rJ   c                 ó:   — g | ]}|                      ¦   «         °|‘ŒS r   )Úis_declaration)ra   Úfns     r   rd   z(HIPBackend.make_llir.<locals>.<listcomp>`  s)   € ÐPÐPÐP�b¸B×<MÒ<MÑ<OÔ<OÐPˆrÐPÐPÐPr   zamdgpu-flat-work-group-sizez1,zamdgpu-waves-per-euzpreserve-signr8   zdenormal-fp-math-f32rM   z
asanrtl.bczocml.bczockl.bcc                 óD   •— g | ]\  }}t          j        ‰|¦  «        ¯|‘ŒS r   )r	   Úneed_extern_lib)ra   rb   ÚpathÚllvm_mods      €r   rd   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
ttg.sharedrœ   )Br   rÆ   rÇ   rÈ   r	   r   rÛ   Úadd_optimize_lds_usager#   ÚconvertÚadd_scf_to_cfÚadd_index_to_llvmirÚadd_allocate_shared_memoryÚadd_to_llvmirrÉ   rÎ   rÑ   Úadd_cf_to_llvmirÚadd_arith_to_llvmirrÓ   rD   ré   Úlower_instruction_sched_hintsr/   r
   ÚcompilationÚdisable_line_infoÚllvmirÚadd_di_scopeÚadd_builtin_func_to_llvmirrÕ   r   Úinit_targetsÚ	to_moduleÚattach_target_tripleÚenable_asanÚattach_datalayoutÚTARGET_TRIPLEÚset_isa_versionÚset_abi_versionÚset_bool_control_constantrK   Úget_functionsÚset_calling_convÚCALLING_CONV_AMDGPU_KERNELÚadd_fn_attrr,   r-   r?   Úadd_fn_target_featureÚadd_fn_asan_attrÚset_all_fn_arg_inregr   rS   rT   rV   Úlink_extern_libsr1   Úoptimize_moduleÚOPTIMIZE_O3r;   Úscalarize_packed_fopsÚ#add_scalarize_packed_fops_llvm_passÚget_int_attrÚcleanup_bitcode_metadataÚdisable_print_inline)rø   r�   rƒ   rÖ   r×   Úcustom_lds_sizeÚ_HIPBackend__HIP_FTZrÇ   Útarget_featuresÚfnsÚdenormal_moder[   Úpathsr  s                @r   Ú	make_llirzHIPBackend.make_llir#  s!  ø€ àˆåŒ_˜Sœ[Ñ)Ô)ˆØ
�ŠÑÔÐð ˆÝŒ
Ô×1Ò1°"°g´lÀOÑTÔTÐTÝŒ×$Ò$ RÑ(Ô(Ð(ÝŒ×*Ò*¨2Ñ.Ô.Ð.åŒ×1Ò1°"Ñ5Ô5Ð5ð ˆ	ÝŒ
Ô×(Ò(¨¨W¬\¸9ÑEÔEÐEÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò˜bÑ!Ô!Ð!åŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×*Ò*¨2Ñ.Ô.Ð.ÝŒ×'Ò'¨Ñ+Ô+Ð+ÝŒ×Ò˜bÑ!Ô!Ð!ÝŒ×$Ò$ RÑ(Ô(Ð(ØÔ ×&Ò&Ñ(Ô(¨FÒ2Ð2ÝŒJÔ×<Ò<¸RÀÄÈwÔOaÑbÔbÐbÝÔ Ô2ð 	+ÝŒM×&Ò& rÑ*Ô*Ð*ÝŒ
Ô×5Ò5°b¸)ÑDÔDÐDØ
�Šˆs‰Œˆõ 	ÔÑÔÐÝ”,‘.”.ˆÝ”> # 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Œ×Ò¥Ô >Ñ?Ô?Ð?ØˆAŒ×ÒÐ8Ð:d¸wÔ?PÐQXÔQbÑ?bÐ:dÐ:dÑeÔeÐeð 	ˆAŒ×ÒÐ0°WÔ5IÐ2KÑLÔLÐLØ+2Ô+EÐQ˜˜È6ˆØˆAŒ×ÒÐ1°=ÑAÔAÐAÝÔÔ(ð 	&Ø�ŒF×(Ò(¨Ñ2Ô2Ð2Ø�ŒF×#Ò#Ñ%Ô%Ð%õ
 	Ô   Q¤Ñ(Ô(Ð(åÔÔ(ð 
	3Ý!¥(™^œ^Ô2°UÑ:ˆNå�N \Ñ1Ñ2Ô2Ý�N YÑ.Ñ/Ô/Ý�N YÑ.Ñ/Ô/ðˆEõ
 Ô! (¨EÑ2Ô2Ð2Ð2ØÔ ð 	3ØiÐiÐiÐi¨gÔ.AÐiÑiÔiˆEÝÔ! (¨EÑ2Ô2Ð2åÔ˜X¥tÔ'7¸¼ÀrÈ2ÈwÔOgÑhÔhÐhåŒ9Ô*ð 	<ÝÔ3°C¸´FÑ;Ô;Ð;ð !×-Ò-¨lÑ;Ô;ˆ�ÑåÔ$ XÑ.Ô.Ð.õ 	Ô  Ñ*Ô*Ð*Ý�8‰}Œ}Ðr   c           	      ón  — t          j        d| ¦  «        }t          |¦  «        dk    sJ ‚|d         |d<   g }|j        dk    r|                     d¦  «         t          j        | t          j        |j	        d||j
        d¦  «        }t          j        j        rt          d	¦  «         t          |¦  «         |S )
Nz3define amdgpu_kernel void @([a-zA-Z_][a-zA-Z0-9_]*)r   r   rb   Ú	attentionzsink-insts-to-avoid-spillsrû   Fz!// -----// AMDGCN Dump //----- //)ÚreÚfindallÚlenrD   Úappendr   Útranslate_to_asmr	   r  r#   r;   r
   Údump_amdgcnÚprint)rø   r�   rƒ   ÚnamesÚflagsÚamdgcns         r   Úmake_amdgcnzHIPBackend.make_amdgcn‘  s³   € õ
 ”
ÐQÐSVÑWÔWˆÝ�5‰zŒz˜QŠˆˆˆØ  œ8ˆ�Ñàˆð
 Ô  KÒ/Ð/Ø�LŠLÐ5Ñ6Ô6Ð6ÝÔ& s­CÔ,=¸w¼|ÈRÐQVÐX_ÔXpÐrwÑxÔxˆÝŒ9Ô ð 	ÝÐ5Ñ6Ô6Ð6Ý�&‰MŒMˆMØˆr   c                 ó”  — d}t           j        j        rd}t          j        | |j        |¦  «        }t                               ¦   «         }t          j	        ¦   «         5 }t          j	        ¦   «         5 }t          |j        d¦  «        5 }|                     |¦  «         d d d ¦  «         n# 1 swxY w Y   t          j        |ddd|j        d|j        g¦  «         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 )	Nrû   rü   Úwbz-flavorÚgnuz-sharedz-oÚrb)r
   r  r  r	   Úassemble_amdgcnr#   rt   rÄ   ÚtempfileÚNamedTemporaryFileÚopenrb   ÚwriteÚ
subprocessÚ
check_callÚread)rø   r�   rƒ   r0  rz   Ú	rocm_pathÚtmp_outÚtmp_inÚfd_inÚfd_outr·   s              r   Ú
make_hsacozHIPBackend.make_hsaco§  s*  € àˆÝÔÔ(ð 	'Ø&ˆOÝÔ# C¨¬°ÑGÔGˆå×/Ò/Ñ1Ô1ˆ	ÝÔ(Ñ*Ô*ð 	$¨gÝÔ,Ñ.Ô.ð q°&Ý˜&œ+ tÑ,Ô,ð '°Ø—K’K Ñ&Ô&Ð&ð'ð 'ð 'ñ 'ô 'ð 'ð 'ð 'ð 'ð 'ð 'øøøð 'ð 'ð 'ð 'åÔ% y°)¸UÀIÈvÌ{Ð\`ÐbiÔbnÐ&oÑpÔpÐpðqð qð qñ qô qð qð qð qð qð qð qøøøð qð qð qð qõ �g”l DÑ)Ô)ð $¨VØ—k’k‘m”m�ð$ð $ð $ñ $ô $ð $ð $ð $ð $ð $ð $øøøð $ð $ð $ð $ð	$ð 	$ð 	$ñ 	$ô 	$ð 	$ð 	$ð 	$ð 	$ð 	$ð 	$øøøð 	$ð 	$ð 	$ð 	$ð ˆ
s~   ÁD=Á1C%ÂB)ÂC%Â)B-Â-C%Â0B-Â1(C%ÃD=Ã%C)	Ã)D=Ã,C)	Ã-D=ÄD&ÄD=Ä&D*	Ä*D=Ä-D*	Ä.D=Ä=EÅEc                 ó¸   ‡ ‡— |t           j        k    rˆˆ fd„|d<   ˆˆ fd„|d<   n|t           j        k    r	ˆˆ fd„|d<   ˆˆ fd„|d<   ˆˆ fd„|d	<   ˆˆ fd
„|d<   d S )Nc                 ó2   •— ‰                      | |‰¦  «        S rŒ   )rØ   ©rø   r�   rƒ   rY   s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>º  s   ø€ °4·>²>À#ÀxÐQXÑ3YÔ3Y€ r   rË   c                 ó2   •— ‰                      | |‰¦  «        S rŒ   )ró   rV  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>»  s   ø€ °D·O²OÀCÈÐSZÑ4[Ô4[€ r   Úttgirc                 ó2   •— ‰                      | |‰¦  «        S rŒ   )rù   rV  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>½  s   ø€ °D·N²NÀ3ÈÐRYÑ4ZÔ4Z€ r   c                 ó2   •— ‰                      | |‰¦  «        S rŒ   )r4  rV  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>¾  s   ø€ ¨t¯~ª~¸cÀ8ÈWÑ/UÔ/U€ r   Úllirc                 ó2   •— ‰                      | |‰¦  «        S rŒ   )rA  rV  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>¿  s   ø€ °×1AÒ1AÀ#ÀxÐQXÑ1YÔ1Y€ r   r@  c                 ó2   •— ‰                      | |‰¦  «        S rŒ   )rS  rV  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>À  s   ø€ °·²ÀÀXÈwÑ0WÔ0W€ r   rz   )r   ÚTRITONÚGLUON)rY   Ústagesrƒ   r•   s   ` ` r   Ú
add_stageszHIPBackend.add_stages¸  s£   øø€ Ø•x”Ò&Ð&ØYÐYÐYÐYÐYˆF�6‰NØ[Ð[Ð[Ð[Ð[ˆF�7‰OˆOØ�œÒ'Ð'ØZÐZÐZÐZÐZˆF�7‰OØUÐUÐUÐUÐUˆˆv‰ØYÐYÐYÐYÐYˆˆxÑØWÐWÐWÐWÐWˆˆw‰ˆˆr   c                 óx   — t          j        t                               ¦   «         dgd¬¦  «        }|› d| j        › �S )Nz	--versionre   )Úencodingr`   )rK  Úcheck_outputrt   rÄ   r   )rY   Úversions     r   rm   zHIPBackend.hashÂ  s?   € åÔ)­:×+FÒ+FÑ+HÔ+HÈ+Ð*VÐahÐiÑiÔiˆØÐ)Ð)˜DœKÐ)Ð)Ð)r   ) rn   ro   rp   Ústaticmethodr   rw   r|   rV   r„   r   rš   rž   r¡   r   r   r¥   r§   r²   rµ   rº   rÄ   rØ   ró   rù   r4  rA  rS  ra  Ú	functoolsÚ	lru_cacherm   Ú__classcell__)r   s   @r   rt   rt   Z   sF  ø€ € € € € àð' 	ð 'ð 'ð 'ñ „\ð'ð"˜yð "¨Tð "ð "ð "ð "ð "ð "ð
%¨#ð %ð %ð %ð %ð" Sð "ð "ð "ð "ð2
ð 
ð 
ð?ð ?ð ?ð>  S¨* _Ô 5ð >ð >ð >ð >ð
ð ð ð ðð ñ „\ðð ðð ñ „\ðð ðð ñ „\ðð ðsð sñ „\ðsð& ðð ñ „\ðð  ð;ð ;ñ „\ð;ðz ðð ñ „\ðð ðkð kñ „\ðkðZ ðð ñ „\ðð* ðð ñ „\ðð Xð Xð Xð €YÔÑÔð*ð *ñ Ôð*ð *ð *ð *ð *r   rt   )Útriton.backends.compilerr   r   r   Útriton._C.libtritonr   r   r   r	   Útritonr
   Údataclassesr   Útypingr   r   r   Útypesr   rh   rG  r7  rK  rg  Úpathlibr   r   r$   r'   r*   rt   r   r   r   ú<module>rq     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ðhð hð hðrð rð rð €�$ÐÑÔð9?ð 9?ð 9?ð 9?ð 9?ñ 9?ô 9?ñ Ôð9?ðxk*ð k*ð k*ð k*ð k*�ñ k*ô k*ð k*ð k*ð k*r   