Ë
    ÚéñiÅY  ã                   óö   — d dl mZmZmZ d dlmZmZmZmZ d dl	m
Z
 d dlmZ d dlmZmZmZ d dlmZ d dlZd dlZd dlZd dlZd dlZd dlmZ d	efd
„Zd„ Zd„ Z ed¬«       G d„ d«      «       Z G d„ de«      Zy)é    )ÚBaseBackendÚ	GPUTargetÚLanguage)ÚirÚpassesÚllvmÚamd)Úknobs)Ú	dataclass)ÚAnyÚDictÚTuple)Ú
ModuleTypeN)ÚPathÚtargetc                 ó   — d„ S )Nc                  ó   — y)N)é   r   r   © )Úlhs_typeÚrhs_types     ú^/var/www/pod-logistic/pod-ai/venv/lib/python3.12/site-packages/triton/backends/amd/compiler.pyú<lambda>z"get_min_dot_size.<locals>.<lambda>   s   � ó    r   ©r   s    r   Úget_min_dot_sizer      s
   € ñ 0Ð/r   c                 óŽ   — t         j                  j                  €| dk(  xs | dk(  xr |du S t         j                  j                  S )NÚgfx942Úgfx950T)r
   r	   Úuse_block_pingpong)ÚarchÚuse_async_copys     r   Úis_pingpong_schedule_enabledr#      sI   € ä—‘×-Ñ-Ð5ð �HÑÒM ¨Ñ!1Ò!L°nÈÐ6Lð XÜ;@¿9¹9×;WÑ;WðXr   c                 ót   — t         j                  j                  €| dk(  S t         j                  j                  S )Nr   )r
   r	   Úuse_in_thread_transpose©r!   s    r   Úis_in_thread_transpose_enabledr'      s.   € Ü!&§¡×!BÑ!BÐ!JˆD�HÑÐqÔPU×PYÑPY×PqÑPqÐqr   T)Úfrozenc                   óR  — e Zd ZU dZeed<   dZeed<   dZeed<   dZeed<   d	Z	e
ed
<   dZeed<   dZeed<   d	Zeed<   dZee   ed<   dZee   ed<   dZeed<   dZee   ed<   dZeed<   dZeed<   dZeed<   dZeed<   dZeed<   dZeed<   dZeed<   d Zeed!<   d"Zeed#<   d$„ Zd%„ Zy	)&Ú
HIPOptionsé   Ú	num_warpsr   Úwaves_per_eué   Ú
num_stagesr   Únum_ctasNÚextern_libsFÚdebugTÚsanitize_overflowr!   )Úfp8e4nvÚfp8e5Úfp8e5b16Úfp8e4b8Úsupported_fp8_dtypesr   Ú!deprecated_fp8_dot_operand_dtypesÚieeeÚdefault_dot_input_precision)r:   Úbf16x3Úbf16x6Úallowed_dot_input_precisionsÚenable_fp_fusionÚlaunch_cooperative_gridÚmatrix_instr_nonkdimÚkpackÚallow_flush_denormÚmax_num_imprecise_acc_defaultÚhipÚbackend_nameÚ Úinstrumentation_modeÚnoneÚschedule_hintc                 óž  — t        | j                  dd «      }|dk\  rdnd}t        j                  | d|«       | j                  dkD  r| j                  | j                  dz
  z  dk(  sJ d	«       ‚| j                  d
k(  rI| j
                  dk7  r:t        j                  d| j
                  › d�«       t        j                  | dd«       t        t        «      j                  dz  }| j                  €i nt        | j                  «      }dD ]  }t        ||› d�z  «      ||<   Œ t        j                  | dt        |j                  «       «      «       y )Né   éþÿÿÿé
   é    é@   Ú	warp_sizer   r   znum_warps must be a power of 2r   zckpack is deprecated starting from gfx950 and will be removed in later releases. So for now kpack = z7 will be overwritten to 1 to make transitioning easier.rB   Úlib)ÚocmlÚocklz.bcr1   )Úintr!   ÚobjectÚ__setattr__r,   rB   ÚwarningsÚwarnr   Ú__file__Úparentr1   ÚdictÚstrÚtupleÚitems)ÚselfÚ	gfx_majorrQ   Údefault_libdirr1   rR   s         r   Ú__post_init__zHIPOptions.__post_init__N   sG  € Ü˜Ÿ	™	 ! B˜Ó(ˆ	Ø# rš/‘B¨rˆ	Ü×Ñ˜4 ¨iÔ8Ø�~‰~ Ò! t§~¡~¸¿¹È!Ñ9KÑ'LÐQRÒ&Rð 	-Ø,ó	-ÐRð �I‰I˜Ò!¨¯
©
°aªÜ�M‰MØuÐvz÷  wAñ  wAð  vBð  Byð  zôô ×Ñ˜t W¨aÔ0äœh›×.Ñ.°Ñ6ˆØ ×,Ñ,Ð4‘b¼$¸t×?OÑ?OÓ:PˆØ#ò 	AˆCÜ" >°s°e¸3°KÑ#?Ó@ˆK˜Òð	Aä×Ñ˜4 ´°k×6GÑ6GÓ6IÓ0JÕKr   c           	      óô   — dj                  | j                  j                  «       D ��cg c]  \  }}|› d|› �‘Œ c}}«      }t        j                  |j                  d«      «      j                  «       S c c}}w )NÚ_ú-úutf-8)ÚjoinÚ__dict__r_   ÚhashlibÚsha256ÚencodeÚ	hexdigest)r`   ÚnameÚvalÚkeys       r   ÚhashzHIPOptions.hasha   s]   € Ø�h‰h¸¿¹×9LÑ9LÓ9N×O©I¨D°#˜4˜&  # šÓOÓPˆÜ�~‰~˜cŸj™j¨Ó1Ó2×<Ñ<Ó>Ð>ùó Ps   ©A4
) Ú__name__Ú
__module__Ú__qualname__r,   rU   Ú__annotations__r-   r/   r0   r1   r\   r2   Úboolr3   r!   r]   r8   r   r9   r;   r>   r?   r@   rA   rB   rC   rD   rF   rH   rJ   rc   rq   r   r   r   r*   r*      sú   … à€IˆsÓØ€L�#ÓØ€J�ÓØ€HˆcÓØ€K�ÓØ€Eˆ4ÓØ"Ð�tÓ"Ø€Dˆ#Óð
 (SÐ˜% ™*ÓRØ46Ð% u¨S¡zÓ6Ø'-Ð Ó-Ø/KÐ  %¨¡*ÓKØ!Ð�dÓ!Ø$)Ð˜TÓ)Ø !Ð˜#Ó!Ø€Eˆ3ƒNØ$Ð˜Ó$Ø)*Ð! 3Ó*Ø€L�#ÓØ "Ð˜#Ó"ð*  €M�3ÓòLó&?r   r*   c                   óJ  ‡ — e Zd ZdZdZedefd„«       Zdeddfˆ fd„Zde	fd„Z
defd„Zd	„ Zd
„ Zdee	ef   fd„Zd„ Zed„ «       Zed„ «       Zed„ «       Zed„ «       Zed„ «       Zed„ «       Zed„ «       Zed„ «       Zed„ «       Zd„ Z ej<                  «       d„ «       Zˆ xZ S )Ú
HIPBackendNFr   c                 ó    — | j                   dk(  S )NrE   )Úbackendr   s    r   Úsupports_targetzHIPBackend.supports_targetj   s   € à�~‰~ Ñ&Ð&r   Úreturnc                 ój   •— t         ‰| �  |«       t        |j                  t        «      sJ ‚d| _        y )NÚhsaco)ÚsuperÚ__init__Ú
isinstancer!   r]   Ú
binary_ext)r`   r   Ú	__class__s     €r   r€   zHIPBackend.__init__n   s+   ø€ Ü‰Ñ˜Ô Ü˜&Ÿ+™+¤sÔ+Ð+Ð+Ø!ˆ�r   c                 ó    — d|j                   › �S )Núhip:r&   ©r`   Úoptionss     r   Úget_target_namezHIPBackend.get_target_names   s   € Ø�g—l‘l�^Ð$Ð$r   c                 óà  — dt         j                  j                  xs | j                  j                  i}|j                  dd«      dkD  rKt        j                  | j                  j                  «      s"t        d| j                  j                  › �«      ‚| j                  j                  dk(  rBt        t        j                  «      }|j                  dh«       t        t        |«      «      |d<   d|vr%t        t        t        j                  «      «      |d<   | j                  j                  d	k(  rCt        t        j                   «      }|j                  d
dh«       t        t        |«      «      |d<   d|vrt         j"                  j$                  |d<   |j                  t        j&                  j)                  «       D �ci c]  }||v sŒ||   €Œ|||   “Œ c}«       t        di |¤ŽS c c}w )Nr!   r0   r   znum_ctas > 1 not supported on r   Útf32r>   r8   r   r6   r7   r9   r?   r   )r
   ÚruntimeÚoverride_archr   r!   Úgetr	   Úsupports_multi_cta_launchÚ
ValueErrorÚsetr*   r>   Úupdater^   Úsortedr8   r9   ÚlanguageÚdefault_fp_fusionÚ__dataclass_fields__Úkeys)r`   ÚoptsÚargsr>   r9   Úks         r   Úparse_optionszHIPBackend.parse_optionsv   s›  € ØœŸ™×3Ñ3ÒG°t·{±{×7GÑ7GÐHˆà�8‰8�J Ó" QÒ&¬s×/LÑ/LÈTÏ[É[×M]ÑM]Ô/^ÜÐ=¸d¿k¹k×>NÑ>NÐ=OÐPÓQÐQð �;‰;×Ñ˜xÒ'Ü+.¬z×/VÑ/VÓ+WÐ(Ø(×/Ñ/°°Ô9Ü38¼Ð@\Ó9]Ó3^ˆDÐ/Ñ0à!¨Ñ-Ü+0´¼
×8WÑ8WÓ1XÓ+YˆDÐ'Ñ(à�;‰;×Ñ˜xÒ'Ü03´J×4`Ñ4`Ó0aÐ-Ø-×4Ñ4°jÀ)Ð5LÔMÜ8=¼fÐEfÓ>gÓ8hˆDÐ4Ñ5à TÑ)Ü',§~¡~×'GÑ'GˆDÐ#Ñ$Ø�‰¬×)HÑ)H×)MÑ)MÓ)OÖu AÐSTÐX\ÒS\ÐaeÐfgÑahÑat�Q˜˜Q™‘ZÒuÔvÜÑ!˜DÑ!Ð!ùò vs   Ç	G+ÇG+ÇG+c                 óH   — |j                   |j                  |j                  fS ©N)r,   r0   Úshared)r`   Úmetadatas     r   Úpack_metadatazHIPBackend.pack_metadata�   s%   € à×ÑØ×ÑØ�O‰Oð
ð 	
r   c                 ó0   — dt        | j                  «      iS )NÚmin_dot_size)r   r   r†   s     r   Úget_codegen_implementationz%HIPBackend.get_codegen_implementation–   s   € ØÔ 0°·±Ó =Ð>Ð>r   c                 ó   — ddl m} d|iS )Nr   )Ú	libdeviceztriton.language.extra.libdevice)Útriton.language.extra.hipr¤   )r`   r¤   s     r   Úget_module_mapzHIPBackend.get_module_map™   s   € Ý7à1°9Ð=Ð=r   c                 óŽ   — t        j                  |«       t        j                  r t        j                  j                  |«       y y rœ   )r	   Úload_dialectsrx   Úinstrumentation)r`   Úctxs     r   r¨   zHIPBackend.load_dialectsž   s2   € Ü×Ñ˜#ÔÜ×%Ò%Ü×&Ñ&×4Ñ4°SÕ9ð &r   c                 óÔ   — dd l }d}t        | d«      r| j                  «       |k  S t        | |j                  «      r-t        | d«      r!| j                  «       j                  «       |k  S y)Nr   iÿÿÿÚ	ptr_rangeÚuntyped_storageF)ÚtorchÚhasattrr¬   r�   ÚTensorr­   Úsize)Úargr®   Ú
MAX_INT_32s      r   Úis_within_2gbzHIPBackend.is_within_2gb£   s]   € ãàˆ
Ü�3˜Ô$Ø—=‘=“? jÑ0Ð0Ü�c˜5Ÿ<™<Ô(¬W°SÐ:KÔ-LØ×&Ñ&Ó(×-Ñ-Ó/°:Ñ=Ð=Ør   c                 óH   — t        j                  | «      }d| v r|ddggz  }|S )NÚSztt.pointer_rangerO   )r   Ú
parse_attr)ÚdescÚrets     r   r·   zHIPBackend.parse_attr®   s1   € ä×$Ñ$ TÓ*ˆØ�$‰;ØÐ'¨Ð,Ð-Ñ-ˆCØˆ
r   c                 óš   — t        j                  | fi |¤Ž}t        j                  j                  rt
        j                  | «      r|dz  }|S )Nr¶   )r   Úget_tensor_specializationr
   r	   Úuse_buffer_opsrx   r´   )r²   Úkwargsr¹   s      r   r»   z$HIPBackend.get_tensor_specializationµ   s?   € ä×3Ñ3°CÑB¸6ÑBˆÜ�9‰9×#Ò#¬
×(@Ñ(@ÀÔ(EØ�3‰JˆCØˆ
r   c                 óô  — t        j                  | j                  «      }|j                  «        t        j
                  j                  |«       t        j                  j                  |«       t        j                  j                  |«       t        j
                  j                  |«       t        j                  j                  |«       t        j                  j                  |«       t        j
                  j                  |«       t        j                  j                  |«       t        j
                  j                  |«       t        j                  j!                  |«       |j#                  | d«       | S )NÚ	make_ttir)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   r¿   zHIPBackend.make_ttir¼   så   € ä�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�‰×!Ñ! "Ô%Ü�‰×.Ñ.¨rÔ2Ü�‰×<Ñ<¸RÔ@Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ Ô#Ü�‰×)Ñ)¨"Ô-Ü�‰×Ñ˜bÔ!Ü�‰×#Ñ# BÔ'Ü�‰×$Ñ$ RÔ(Ü�‰×#Ñ# BÔ'Ø
�‰ˆs�KÔ Øˆ
r   c                 óä  — t        j                  | j                  «      }|j                  «        t        j
                  j                  |d|j                  › �|j                  |j                  |j                  «       |j                  | d«       t        j                  | j                  «      }|j                  «        d}t        j                  j                  |«       t        j                  j                  ||«       t        j                  j                  |«       t        j                  j!                  |«       t"        j                  j                  j%                  ||j                  |j&                  |j(                  «       t        j                  j                  |«       t"        j                  j                  j+                  |«       t"        j                  j                  j-                  ||j                  «       t"        j                  j                  j/                  |«       t        j                  j1                  |«       t        j2                  j5                  |«       t        j
                  j7                  |«       t        j2                  j5                  |«       t8        j"                  j:                  }t=        |j                  |«      }t"        j                  j                  j?                  ||j@                  «       t"        j                  j                  jC                  |||«       |r4t"        j                  j                  jE                  ||j                  «       t        j2                  j5                  |«       |jF                  jI                  «       dk7  rJ|jF                  jK                  d«      D ],  }t"        j                  j                  jM                  ||«       Œ. t        j                  j                  |«       t        j                  jO                  |«       tQ        |j                  «      rHt"        j                  j                  jS                  |«       t        j                  j                  |«       t"        j                  j                  jU                  |«       |rC|j@                  dkD  r4t"        j                  j                  jW                  ||j@                  «       t8        j"                  jX                  r®t"        j                  j                  j[                  |«       t        j2                  j5                  |«       t"        j                  j                  j]                  ||j                  t8        j"                  j^                  t8        j"                  j`                  «       t"        j                  j                  jc                  |«       t        j2                  j5                  |«       t        j2                  je                  |«       t        j2                  jg                  |«       |j                  | d«       | ji                  «       |d<   | S )	Nr…   Úmake_ttgir_earlyFrI   ú,r   Ú
make_ttgirÚtensordesc_meta)5r   rÀ   rÁ   rÂ   r   rÅ   Úadd_convert_to_ttgpuirr!   r,   rQ   r0   rÏ   ÚttgpuirÚadd_coalesceÚadd_f32_dot_tcÚadd_remove_layout_conversionsÚadd_optimize_thread_localityr	   Úadd_accelerate_matmulrA   rB   Úadd_optimize_epilogueÚadd_optimize_dot_operandsÚadd_hoist_layout_conversionsÚadd_fuse_nested_loopsrÃ   rÈ   rÌ   r
   r"   r#   Úadd_schedule_loopsr/   Úadd_pipelineÚadd_coalesce_async_copyrJ   ÚlowerÚsplitÚinsert_instruction_sched_hintsÚadd_reduce_data_duplicationr'   Úadd_in_thread_transposeÚadd_reorder_instructionsÚadd_block_pingpongr¼   Úadd_canonicalize_pointersÚadd_convert_to_buffer_opsÚuse_buffer_atomicsÚ%buffer_ops_analyze_small_tensor_rangeÚadd_fold_true_cmpirË   rÍ   Úget_tensordesc_metadata)rÐ   rž   r‡   rÑ   ÚemuTF32r"   r    Úhints           r   rÕ   zHIPBackend.make_ttgirÍ   s  € ä�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�‰×*Ñ*¨2°°g·l±l°^Ð/DÀg×FWÑFWÐY`×YjÑYjØ+2×+;Ñ+;ô	=à
�‰ˆsÐ&Ô'Ü�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔØˆÜ�‰×#Ñ# BÔ'Ü�‰×%Ñ% b¨'Ô2Ü�‰×4Ñ4°RÔ8Ü�‰×3Ñ3°BÔ7Ü�
‰
×Ñ×0Ñ0°°W·\±\À7×C_ÑC_Ðah×anÑanÔoÜ�‰×4Ñ4°RÔ8Ü�
‰
×Ñ×0Ñ0°Ô4Ü�
‰
×Ñ×4Ñ4°R¸¿¹ÔFÜ�
‰
×Ñ×7Ñ7¸Ô;ä�‰×,Ñ,¨RÔ0Ü�‰×'Ñ'¨Ô+Ü�‰×#Ñ# BÔ'Ü�‰×'Ñ'¨Ô+äŸ™×1Ñ1ˆÜ9¸'¿,¹,ÈÓWÐä�
‰
×Ñ×-Ñ-¨b°'×2DÑ2DÔEÜ�
‰
×Ñ×'Ñ'¨¨NÐ<NÔOÙÜ�J‰J×Ñ×6Ñ6°r¸7¿<¹<ÔHÜ�‰×'Ñ'¨Ô+Ø× Ñ ×&Ñ&Ó(¨FÒ2Ø×-Ñ-×3Ñ3°CÓ8ò L�Ü—
‘
×"Ñ"×AÑAÀ"ÀdÕKðLä�‰×4Ñ4°RÔ8Ü�‰×2Ñ2°2Ô6Ü)¨'¯,©,Ô7Ü�J‰J×Ñ×6Ñ6°rÔ:Ü�N‰N×8Ñ8¸Ô<Ü�
‰
×Ñ×3Ñ3°BÔ7Ù '×"4Ñ"4°qÒ"8Ü�J‰J×Ñ×1Ñ1°"°g×6HÑ6HÔIä�9‰9×#Ò#Ü�J‰J×Ñ×8Ñ8¸Ô<Ü�M‰M×+Ñ+¨BÔ/Ü�J‰J×Ñ×8Ñ8ØØ—‘Ü—	‘	×,Ñ,Ü—	‘	×?Ñ?ô	ô 	�
‰
×Ñ×-Ñ-¨bÔ1Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ø
�‰ˆs�LÔ!Ø&)×&AÑ&AÓ&CˆÐ"Ñ#Øˆ
r   c                 ó&  — | }t        j                  |j                  «      }|j                  «        t        j
                  j                  |«       t        j
                  j                  |«       t        j                  j                  |«       t        j                  j                  |«       t        j
                  j                  |«       t        j                  j                  |«       |j                  |d«       |j!                  «       |d<   |S )NÚgluon_to_ttgirrÖ   )r   rÀ   rÁ   rÂ   r   ÚgluonrÄ   Úadd_resolve_auto_encodingsrÃ   Úadd_sccprÅ   Úadd_loop_aware_cserÈ   rØ   Ú add_combine_tensor_select_and_ifrÏ   rñ   )Úsrcrž   r‡   rÐ   rÑ   s        r   rõ   zHIPBackend.gluon_to_ttgir  s·   € àˆÜ�_‰_˜SŸ[™[Ó)ˆØ
�‰Ôä�‰× Ñ  Ô$Ü�‰×/Ñ/°Ô3Ü�‰×Ñ˜rÔ"Ü�‰×&Ñ& rÔ*Ü�‰×&Ñ& rÔ*Ü�‰×7Ñ7¸Ô;à
�‰ˆsÐ$Ô%Ø&)×&AÑ&AÓ&CˆÐ"Ñ#Øˆ
r   c                 óÄ  — | }t        j                  |j                  «      }|j                  «        t        j
                  j                  j                  ||j                  «       d}t        j
                  j                  j                  ||j                  |«       t
        j                  j                  |«       t
        j                  j                  |«       t
        j                  j                  |«       t        j
                  j                  j                  |«       t         j"                  r+t         j"                  j%                  d||j                  «       d}t        j
                  j                  j'                  ||j                  |«       t
        j(                  j+                  |«       t
        j(                  j-                  |«       t
        j                  j/                  |«       t
        j                  j1                  |«       t
        j(                  j+                  |«       t
        j(                  j-                  |«       t
        j(                  j3                  |«       |j4                  j7                  «       dk7  r?t        j
                  j                  j9                  ||j                  |j:                  «       t         j"                  r+t         j"                  j%                  d||j                  «       t<        j>                  j@                  s9t<        j>                  jB                  st
        jD                  jG                  |«       t        j
                  j                  jI                  ||«       |jK                  |d«       t<        j>                  jB                  rÚt<        j>                  j@                  s`t        j                  |j                  «      }|j                  «        t
        jD                  jG                  |«       |jK                  |d«       t        j                  |j                  «      }|j                  «        t
        jD                  jM                  |«       |jK                  |d«       tO        jP                  «        tO        j                  «       }tO        jR                  ||«      }t	        jT                  |«       d	}	t<        j>                  jV                  rd
}	tO        jX                  |t        jZ                  |j                  |	«       t	        j\                  ||j                  «       t	        j^                  |d«       t	        j`                  |dd«       t	        j`                  |dd«       t	        j`                  |dd«       t	        j`                  |d|jb                  dk(  «       |je                  «       D �
cg c]  }
|
jg                  «       rŒ|
‘Œ }}
|d   ji                  t        jj                  «       |d   jm                  dd|jn                  |jb                  z  › �«       d|j4                  jq                  d«      v r|d   jm                  dd«       |d   jm                  dd«       |d   jm                  d|jr                  › d|jr                  › �«       |jt                  rdnd}|d   jm                  d|«       t<        j>                  jV                  r'|d   jw                  d
«       |d   jy                  «        t	        jz                  |d   «       t<        j>                  jV                  r\t}        t~        «      j€                  dz  }tƒ        |d z  «      tƒ        |d!z  «      tƒ        |d"z  «      g}tO        j„                  ||«       ne|j†                  rY|j†                  D ��cg c]  \  }}t	        jˆ                  ||«      sŒ|‘Œ  }}}t‹        |«      dkD  rtO        j„                  ||«       tO        jŒ                  |tN        jŽ                  |j                  d	g |j�                  «       t	        j’                  |j                  «      r<|d   j•                  d#«       |d   j•                  d$«       |d   j•                  d%«       t<        j                  j–                  rt	        j˜                  |d   «       | j›                  d&«      |d'<   | j›                  d(«      xs d|d)<   | j›                  d*«      xs d+|d,<   t	        jœ                  |«       t	        jž                  |«       tƒ        |«      S c c}
w c c}}w )-Nr   Úttgpuir_to_llvmirTrI   Úllvmir_to_llvmÚ	make_llirzmake_llir.disable_line_infoz,make_llir.dump_ir_extract_di_local_variablesrG   ú+xnackiô  Ú__oclc_finite_only_optFÚ__oclc_correctly_rounded_sqrt32Ú__oclc_unsafe_math_optÚ__oclc_wavefrontsize64rP   z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-f32rR   z
asanrtl.bczocml.bczockl.bczamdgpu-no-workgroup-id-xzamdgpu-no-workgroup-id-yzamdgpu-no-workgroup-id-zz
ttg.sharedr�   zttg.profile_scratch_memory_sizeÚprofile_scratch_sizez$ttg.profile_scratch_memory_alignmentr   Úprofile_scratch_align)Pr   rÀ   rÁ   rÂ   r	   r   rØ   Úadd_update_async_wait_countr!   Úadd_optimize_lds_usageÚconvertÚadd_scf_to_cfrö   rÄ   Úadd_index_to_llvmirÚadd_allocate_shared_memoryrx   r©   ÚpatchÚadd_to_llvmirrÃ   rÈ   rË   Úadd_cf_to_llvmirÚadd_arith_to_llvmirrÍ   rJ   rå   Úlower_instruction_sched_hintsr/   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_constantrQ   Úget_functionsÚis_declarationÚset_calling_convÚCALLING_CONV_AMDGPU_KERNELÚadd_fn_attrr,   ræ   r-   rC   Úadd_fn_target_featureÚadd_fn_asan_attrÚset_all_fn_arg_inregr   rZ   r[   r]   Úlink_extern_libsr1   Úneed_extern_libÚlenÚoptimize_moduleÚOPTIMIZE_O3r?   Úhas_architected_sgprsÚremove_fn_attrÚscalarize_packed_fopsÚ#add_scalarize_packed_fops_llvm_passÚget_int_attrÚcleanup_bitcode_metadataÚdisable_print_inline)rû   rž   r‡   rÐ   rÑ   Úcustom_lds_sizeÚ_HIPBackend__HIP_FTZrÁ   Úllvm_modÚtarget_featuresÚfnÚfnsÚdenormal_moderb   Úpathsrn   Úpaths                    r   rÿ   zHIPBackend.make_llir  sà  € àˆä�_‰_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�
‰
×Ñ×6Ñ6°r¸7¿<¹<ÔHð ˆÜ�
‰
×Ñ×1Ñ1°"°g·l±lÀOÔTÜ�‰×$Ñ$ RÔ(Ü�‰× Ñ  Ô$Ü�‰×*Ñ*¨2Ô.ä�
‰
×Ñ×5Ñ5°bÔ9ä×%Ò%Ü×&Ñ&×,Ñ,Ð-@À"ÀcÇkÁkÔRð ˆ	Ü�
‰
×Ñ×(Ñ(¨¨W¯\©\¸9ÔEÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!ä�‰×'Ñ'¨Ô+Ü�‰×*Ñ*¨2Ô.Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(à× Ñ ×&Ñ&Ó(¨FÒ2Ü�J‰J×Ñ×<Ñ<¸RÀÇÁÈw×OaÑOaÔbô ×%Ò%Ü×&Ñ&×,Ñ,Ð-=¸rÀ3Ç;Á;ÔOä× Ñ ×2Ò2¼5×;LÑ;L×;oÒ;oÜ�M‰M×&Ñ& rÔ*ä�
‰
×Ñ×5Ñ5°b¸)ÔDØ
�‰ˆs�KÔ ä×Ñ×?Ò?ä×$Ñ$×6Ò6Ü—_‘_ S§[¡[Ó1�Ø—‘Ô!Ü—‘×*Ñ*¨2Ô.Ø—‘�sÐ9Ô:ô —‘ §¡Ó-ˆBØ�O‰OÔÜ�M‰M×/Ñ/°Ô3Ø�F‰F�3ÐFÔGô 	×ÑÔÜ—,‘,“.ˆÜ—>‘> # wÓ/ˆÜ× Ñ  Ô*ØˆÜ×Ñ×(Ò(Ø&ˆOÜ×Ñ˜x¬×):Ñ):¸G¿L¹LÈ/ÔZô 	×Ñ˜H g§l¡lÔ3Ü×Ñ˜H cÔ*Ü×%Ñ% hÐ0HÈ%ÔPÜ×%Ñ% hÐ0QÐSWÔXÜ×%Ñ% hÐ0HÈ%ÔPÜ×%Ñ% hÐ0HÈ'×J[ÑJ[Ð_aÑJaÔbð %×2Ñ2Ó4ÖP�b¸B×<MÑ<MÕ<OŠrÐPˆÐPàˆA‰×Ñ¤× >Ñ >Ô?ØˆA‰×ÑÐ8¸B¸w×?PÑ?PÐQX×QbÑQbÑ?bÐ>cÐ:dÔeØ# w×'<Ñ'<×'BÑ'BÀ3Ó'GÑGØ�‰F×ÑÐ6¸ÔHØˆA‰×ÑÐ4°fÔ=ð 	ˆA‰×ÑÐ0°W×5IÑ5IÐ4JÈ"ÈW×MaÑMaÐLbÐ2cÔdØ+2×+EÒ+E™È6ˆØˆA‰×ÑÐ1°=ÔAÜ×Ñ×(Ò(Ø�‰F×(Ñ(¨Ô2Ø�‰F×#Ñ#Ô%ô
 	× Ñ   Q¡Ô(ä×Ñ×(Ò(Ü!¤(›^×2Ñ2°UÑ:ˆNä�N \Ñ1Ó2Ü�N YÑ.Ó/Ü�N YÑ.Ó/ðˆEô
 ×!Ñ! (¨EÕ2Ø× Ò Ø.5×.AÑ.A×i™l˜t TÄS×EXÑEXÐYaÐcgÕEh’TÐiˆEÑiÜ�5‹z˜AŠ~Ü×%Ñ% h°Ô6ä×Ñ˜X¤t×'7Ñ'7¸¿¹ÀrÈ2Èw×OgÑOgÔhô ×$Ñ$ W§\¡\Ô2Ø�‰F×!Ñ!Ð"<Ô=Ø�‰F×!Ñ!Ð"<Ô=Ø�‰F×!Ñ!Ð"<Ô=ä�9‰9×*Ò*Ü×3Ñ3°C¸±FÔ;ð !×-Ñ-¨lÓ;ˆ�ÑØ+.×+;Ñ+;Ð<]Ó+^Ò+cÐbcˆÐ'Ñ(Ø,/×,<Ñ,<Ð=cÓ,dÒ,iÐhiˆÐ(Ñ)ä×$Ñ$ XÔ.ô 	× Ñ  Ô*Ü�8‹}ÐùòA QùóJ js   ×e×0eßeß*ec           	      óÒ  — t        j                  d| «      }t        |«      dk(  sJ ‚|d   |d<   g }d|j                  v rdnd}t	        j
                  | j                  d«      «      j                  «       }|d   d	z   |z   }t        j                  | t        j                  |j                  |||j                  |«      }t        j                  | t        j                  |j                  |||j                  |«       t        j                  | 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   rn   Úgfx11z-real-true16rG   rg   re   Fz!// -----// AMDGCN Dump //----- //)ÚreÚfindallr-  r!   rj   rk   rl   rm   r   Útranslate_to_mirr	   r  r?   Údump_sched_dagÚtranslate_to_asmr
   Údump_amdgcnÚprint)
rû   rž   r‡   ÚnamesÚflagsÚfeaturesÚir_hashÚdump_file_idre   Úamdgcns
             r   Úmake_amdgcnzHIPBackend.make_amdgcnº  s/  € ô
 —
‘
ÐQÐSVÓWˆÜ�5‹z˜QŠÐˆØ  ™8ˆ�ÑàˆØ%,°·±Ñ%<‘>À"ˆÜ—.‘. §¡¨GÓ!4Ó5×?Ñ?ÓAˆØ˜Q‘x #‘~¨Ñ/ˆÜ×!Ñ! #¤s×'8Ñ'8¸'¿,¹,ÈÐRWÐY`×YqÑYqØ".ó0ˆä×Ñ˜C¤×!2Ñ!2°G·L±LÀ(ÈEÐSZ×SkÑSkØ(ô	*ä×&Ñ& s¬C×,=Ñ,=¸w¿|¹|ÈXÐW\Ð^e×^vÑ^vØ',ó.ˆä�9‰9× Ò ÜÐ5Ô6Ü�&ŒMØˆr   c                 ól  — d}t         j                  j                  rd}t        j                  | |j
                  |«      }t        j                  «       5 }t        j                  «       5 }t        |j                  d«      5 }|j                  |«       d d d «       t        j                  |j                  |j                  «       d d d «       t        |j                  d«      5 }|j                  «       }	d d d «       d d d «       	S # 1 sw Y   ŒtxY w# 1 sw Y   ŒNxY w# 1 sw Y   Œ+xY w# 1 sw Y   	S xY w)NrG   r   ÚwbÚrb)r
   r  r  r	   Úassemble_amdgcnr!   ÚtempfileÚNamedTemporaryFileÚopenrn   ÚwriteÚ
link_hsacoÚread)
rû   rž   r‡   r:  r~   Útmp_outÚtmp_inÚfd_inÚfd_outr¹   s
             r   Ú
make_hsacozHIPBackend.make_hsacoÒ  sû   € àˆÜ×Ñ×(Ò(Ø&ˆOÜ×#Ñ# C¨¯©°ÓGˆÜ×(Ñ(Ó*ð 	$¨gÜ×,Ñ,Ó.ð :°&Ü˜&Ÿ+™+ tÓ,ð '°Ø—K‘K Ô&÷'ä—‘˜vŸ{™{¨G¯L©LÔ9÷:ô �g—l‘l DÓ)ð $¨VØ—k‘k“m�÷$÷	$ð ˆ
÷'ð 'ú÷:ð :ú÷$ð $ú÷	$ð ˆ
úsT   ÁD)Á)DÂ DÂ2DÃD)Ã"DÃ3D)ÄDÄ
DÄD	ÄD)ÄD&	Ä"D)Ä)D3c                 ó>  ‡ ‡— |t         j                  k(  rˆˆ fd„|d<   ˆˆ fd„|d<   n|t         j                  k(  r	ˆˆ fd„|d<   ˆˆ fd„|d<   ˆˆ fd„|d	<   ˆˆ fd
„|d<   t        j                  j
                  �$t        j                  j                  ‰ |‰|d «       y y )Nc                 ó*   •— ‰j                  | |‰«      S rœ   )r¿   ©rû   rž   r‡   r`   s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>ã  s   ø€ °4·>±>À#ÀxÐQXÓ3Y€ r   rÅ   c                 ó*   •— ‰j                  | |‰«      S rœ   )rÕ   ra  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>ä  s   ø€ °D·O±OÀCÈÐSZÓ4[€ r   Úttgirc                 ó*   •— ‰j                  | |‰«      S rœ   )rõ   ra  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>æ  s   ø€ °D×4GÑ4GÈÈXÐW^Ó4_€ r   c                 ó*   •— ‰j                  | |‰«      S rœ   )rÿ   ra  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>ç  s   ø€ ¨t¯~©~¸cÀ8ÈWÓ/U€ r   Úllirc                 ó*   •— ‰j                  | |‰«      S rœ   )rO  ra  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>è  s   ø€ °×1AÑ1AÀ#ÀxÐQXÓ1Y€ r   rN  c                 ó*   •— ‰j                  | |‰«      S rœ   )r^  ra  s     €€r   r   z'HIPBackend.add_stages.<locals>.<lambda>é  s   ø€ °·±ÀÀXÈwÓ0W€ r   r~   )r   ÚTRITONÚGLUONr
   r‹   Úadd_stages_inspection_hook)r`   Ústagesr‡   r“   s   ` ` r   Ú
add_stageszHIPBackend.add_stagesá  sƒ   ù€ Ø”x—‘Ò&ÜYˆF�6‰NÜ[ˆF�7ŠOØœŸ™Ò'Ü_ˆF�7‰OÜUˆˆv‰ÜYˆˆxÑÜWˆˆw‰Ü�=‰=×3Ñ3Ð?Ü�M‰M×4Ñ4°T¸6À7ÈHÐVZÕ[ð @r   c                 ó   — | j                   › S rœ   r   )r`   s    r   rq   zHIPBackend.hashí  s   € à—+‘+�Ðr   )!rr   rs   rt   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ÿ   rO  r^  rm  Ú	functoolsÚ	lru_cacherq   Ú__classcell__)rƒ   s   @r   rx   rx   f   sS  ø„ Ø€OØ,1Ð)àð' 	ò 'ó ð'ð"˜yð "¨Tõ "ð
%¨#ó %ð" Só "ò2
ò?ð>  S¨* _Ñ 5ó >ò
:ð
 ñó ðð ñó ðð ñó ðð ñó ðð  ñ<ó ð<ð| ñó ðð  ñZó ðZðx ñó ðð. ñó ðò
\ð €Y×ÑÓñ ó ô r   rx   )Útriton.backends.compilerr   r   r   Útriton._C.libtritonr   r   r   r	   Útritonr
   Údataclassesr   Útypingr   r   r   Útypesr   rj   rT  rB  rq  rX   Úpathlibr   r   r#   r'   r*   rx   r   r   r   ú<module>r{     sz   ðß EÑ Eß 5Ó 5Ý Ý !ß #Ñ #Ý Û Û Û 	Û Û Ý ð0˜Yó 0òXò
rñ �$Ô÷D?ð D?ó ðD?ôNI �õ I r   