ó
    �~iÅY  ã                   óò   • S SK JrJrJr  S SKJrJrJrJr  S SK	J
r
  S SKJr  S SKJrJrJr  S SKJr  S SKrS SKrS SKrS SKrS SKrS SKJr  S	\4S
 jrS rS r\" SS9 " S S5      5       r " S S\5      rg)é    )ÚBaseBackendÚ	GPUTargetÚLanguage)ÚirÚpassesÚllvmÚamd)Úknobs)Ú	dataclass)ÚAnyÚDictÚTuple)Ú
ModuleTypeN)ÚPathÚtargetc                 ó   • S $ )Nc                 ó   • g)N)é   r   r   © )Úlhs_typeÚrhs_types     ÚY/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/backends/amd/compiler.pyÚ<lambda>Ú"get_min_dot_size.<locals>.<lambda>   s   €  ió    r   ©r   s    r   Úget_min_dot_sizer      s
   € ñ 0Ð/r   c                 ó¤   • [         R                  R                  c  U S:H  =(       d    U S:H  =(       a    USL $ [         R                  R                  $ )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                 óv   • [         R                  R                  c  U S:H  $ [         R                  R                  $ )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                   óZ  • \ rS rSr% Sr\\S'   Sr\\S'   Sr\\S'   Sr	\\S	'   S
r
\\S'   Sr\\S'   Sr\\S'   S
r\\S'   Sr\\   \S'   Sr\\   \S'   Sr\\S'   Sr\\   \S'   Sr\\S'   Sr\\S'   Sr\\S'   Sr\\S'   Sr\\S'   Sr\\S'   Sr\\S '   S!r\\S"'   S#r\\S$'   S% rS& r Sr!g
)'Ú
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                 ó¬  • [        U R                  SS 5      nUS:¼  a  SOSn[        R                  U SU5        U R                  S:”  a   U R                  U R                  S-
  -  S:X  d   S	5       eU R                  S
:X  aK  U R
                  S:w  a;  [        R                  " SU R
                   S35        [        R                  U SS5        [        [        5      R                  S-  nU R                  c  0 O[        U R                  5      nS H  n[        X5 S3-  5      XE'   M     [        R                  U S[        UR                  5       5      5        g )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.rD   Úlib)ÚocmlÚocklz.bcr3   )Úintr"   ÚobjectÚ__setattr__r.   rD   ÚwarningsÚwarnr   Ú__file__Úparentr3   ÚdictÚstrÚtupleÚitems)ÚselfÚ	gfx_majorrS   Údefault_libdirr3   rT   s         r   Ú__post_init__ÚHIPOptions.__post_init__N   s<  € Ü˜Ÿ	™	 ! 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ˆÛ#ˆCÜ" >°e¸3°KÑ#?Ó@ˆKÓñ $ä×Ñ˜4 ´°k×6GÑ6GÓ6IÓ0JÕKr   c           	      óú   • SR                  U R                  R                  5        VVs/ s H  u  pU SU 3PM     snn5      n[        R                  " UR                  S5      5      R                  5       $ s  snnf )NÚ_Ú-úutf-8)ÚjoinÚ__dict__ra   ÚhashlibÚsha256ÚencodeÚ	hexdigest)rb   ÚnameÚvalÚkeys       r   ÚhashÚHIPOptions.hasha   sa   € Ø�h‰h¸¿¹×9LÑ9LÔ9NÔOÒ9N©I¨D˜4˜&  # ›Ñ9NÒOÓPˆÜ�~Š~˜cŸj™j¨Ó1Ó2×<Ñ<Ó>Ð>ùó Ps   ©A7
)"Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__r.   rW   Ú__annotations__r/   r1   r2   r3   r^   r4   Úboolr5   r"   r_   r:   r   r;   r=   r@   rA   rB   rC   rD   rE   rF   rH   rJ   rL   re   rt   Ú__static_attributes__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                   ó^  ^ • \ rS rSrSrSr\S\4S j5       rS\SS4U 4S jjr	S\
4S jrS\4S	 jrS
 rS rS\\
\4   4S jrS r\S 5       r\S 5       r\S 5       r\S 5       r\S 5       r\S 5       r\S 5       r\S 5       r\S 5       rS r\R>                  " 5       S 5       r Sr!U =r"$ )Ú
HIPBackendéf   NFr   c                 ó    • U R                   S:H  $ )NrG   )Úbackendr   s    r   Úsupports_targetÚHIPBackend.supports_targetj   s   € à�~‰~ Ñ&Ð&r   Úreturnc                 ót   >• [         TU ]  U5        [        UR                  [        5      (       d   eSU l        g )NÚhsaco)ÚsuperÚ__init__Ú
isinstancer"   r_   Ú
binary_ext)rb   r   Ú	__class__s     €r   rˆ   ÚHIPBackend.__init__n   s.   ø€ Ü‰Ñ˜Ô Ü˜&Ÿ+™+¤s×+Ñ+Ð+Ð+Ø!ˆ�r   c                 ó    • SUR                    3$ )Núhip:r'   ©rb   Úoptionss     r   Úget_target_nameÚHIPBackend.get_target_names   s   € Ø�g—l‘l�^Ð$Ð$r   c                 ó  • S[         R                  R                  =(       d    U R                  R                  0nUR                  SS5      S:”  aQ  [        R                  " U R                  R                  5      (       d"  [        SU R                  R                   35      eU R                  R                  S:X  aB  [        [        R                  5      nUR                  S15        [        [        U5      5      US'   SU;  a%  [        [        [        R                  5      5      US'   U R                  R                  S	:X  aC  [        [        R                   5      nUR                  S
S15        [        [        U5      5      US'   SU;  a  [         R"                  R$                  US'   UR                  [        R&                  R)                  5        Vs0 s H  oUU;   d  M
  X   c  M  XQU   _M     sn5        [        S0 UD6$ s  snf )Nr"   r2   r   znum_ctas > 1 not supported on r   Útf32r@   r:   r    r8   r9   r;   rA   r   )r
   ÚruntimeÚoverride_archr   r"   Úgetr	   Úsupports_multi_cta_launchÚ
ValueErrorÚsetr+   r@   Úupdater`   Úsortedr:   r;   ÚlanguageÚdefault_fp_fusionÚ__dataclass_fields__Úkeys)rb   ÚoptsÚargsr@   r;   Úks         r   Úparse_optionsÚ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Ò)O AÐX\ÑS\“ZÐaeÑah“Z�Q˜Q™’ZÑ)OÑuÔvÜÑ!˜DÑ!Ð!ùò vs   Ç	HÇ HÇ'	Hc                 óH   • UR                   UR                  UR                  4$ ©N)r.   r2   Úshared)rb   Úmetadatas     r   Úpack_metadataÚHIPBackend.pack_metadata�   s%   € à×ÑØ×ÑØ�O‰Oð
ð 	
r   c                 ó0   • S[        U R                  5      0$ )NÚmin_dot_size)r   r   r�   s     r   Úget_codegen_implementationÚ%HIPBackend.get_codegen_implementation–   s   € ØÔ 0°·±Ó =Ð>Ð>r   c                 ó   • SSK Jn  SU0$ )Nr   )Ú	libdeviceztriton.language.extra.libdevice)Útriton.language.extra.hipr±   )rb   r±   s     r   Úget_module_mapÚHIPBackend.get_module_map™   s   € Ý7à1°9Ð=Ð=r   c                 óš   • [         R                  " U5        [        R                  (       a   [        R                  R                  U5        g g r§   )r	   Úload_dialectsr~   Úinstrumentation)rb   Úctxs     r   r¶   ÚHIPBackend.load_dialectsž   s2   € Ü×Ò˜#ÔÜ×%×%Ü×&Ñ&×4Ñ4°SÕ9ð &r   c                 óð   • SS K nSn[        U S5      (       a  U R                  5       U:*  $ [        XR                  5      (       a2  [        U S5      (       a!  U R                  5       R                  5       U:*  $ g)Nr   iÿÿÿÚ	ptr_rangeÚuntyped_storageF)ÚtorchÚhasattrr»   r‰   ÚTensorr¼   Úsize)Úargr½   Ú
MAX_INT_32s      r   Úis_within_2gbÚHIPBackend.is_within_2gb£   sd   € ãàˆ
Ü�3˜×$Ñ$Ø—=‘=“? jÑ0Ð0Ü�cŸ<™<×(Ñ(¬W°SÐ:K×-LÑ-LØ×&Ñ&Ó(×-Ñ-Ó/°:Ñ=Ð=Ør   c                 óN   • [         R                  " U 5      nSU ;   a  USS//-  nU$ )NÚSztt.pointer_rangerQ   )r   Ú
parse_attr)ÚdescÚrets     r   rÇ   ÚHIPBackend.parse_attr®   s1   € ä×$Ò$ TÓ*ˆØ�$‹;ØÐ'¨Ð,Ð-Ñ-ˆCØˆ
r   c                 ó°   • [         R                  " U 40 UD6n[        R                  R                  (       a  [
        R                  U 5      (       a  US-  nU$ )NrÆ   )r   Úget_tensor_specializationr
   r	   Úuse_buffer_opsr~   rÃ   )rÁ   ÚkwargsrÉ   s      r   rÌ   Ú$HIPBackend.get_tensor_specializationµ   sB   € ä×3Ò3°CÑB¸6ÑBˆÜ�9‰9×#×#¬
×(@Ñ(@À×(EÑ(EØ�3‰JˆCØˆ
r   c                 óö  • [         R                  " U R                  5      nUR                  5         [        R
                  R                  U5        [        R                  R                  U5        [        R                  R                  U5        [        R
                  R                  U5        [        R                  R                  U5        [        R                  R                  U5        [        R
                  R                  U5        [        R                  R                  U5        [        R
                  R                  U5        [        R                  R!                  U5        UR#                  U S5        U $ )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Ñ   ÚHIPBackend.make_ttir¼   så   € ä�_Š_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�‰×!Ñ! "Ô%Ü�‰×.Ñ.¨rÔ2Ü�‰×<Ñ<¸RÔ@Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ Ô#Ü�‰×)Ñ)¨"Ô-Ü�‰×Ñ˜bÔ!Ü�‰×#Ñ# BÔ'Ü�‰×$Ñ$ RÔ(Ü�‰×#Ñ# BÔ'Ø
�‰ˆs�KÔ Øˆ
r   c                 ó  • [         R                  " U R                  5      nUR                  5         [        R
                  R                  USUR                   3UR                  UR                  UR                  5        UR                  U S5        [         R                  " U R                  5      nUR                  5         Sn[        R                  R                  U5        [        R                  R                  X45        [        R                  R                  U5        [        R                  R!                  U5        ["        R                  R                  R%                  X2R                  UR&                  UR(                  5        [        R                  R                  U5        ["        R                  R                  R+                  U5        ["        R                  R                  R-                  X2R                  5        ["        R                  R                  R/                  U5        [        R                  R1                  U5        [        R2                  R5                  U5        [        R
                  R7                  U5        [        R2                  R5                  U5        [8        R"                  R:                  n[=        UR                  U5      n["        R                  R                  R?                  X2R@                  5        ["        R                  R                  RC                  X5U5        U(       a3  ["        R                  R                  RE                  X2R                  5        [        R2                  R5                  U5        URF                  RI                  5       S:w  aK  URF                  RK                  S5       H,  n["        R                  R                  RM                  X75        M.     [        R                  R                  U5        [        R                  RO                  U5        [Q        UR                  5      (       aH  ["        R                  R                  RS                  U5        [        R                  R                  U5        ["        R                  R                  RU                  U5        U(       aC  UR@                  S:”  a3  ["        R                  R                  RW                  X2R@                  5        [8        R"                  RX                  (       a®  ["        R                  R                  R[                  U5        [        R2                  R5                  U5        ["        R                  R                  R]                  UUR                  [8        R"                  R^                  [8        R"                  R`                  5        ["        R                  R                  Rc                  U5        [        R2                  R5                  U5        [        R2                  Re                  U5        [        R2                  Rg                  U5        UR                  U S5        U Ri                  5       US'   U $ )	NrŽ   Úmake_ttgir_earlyFrK   Ú,r   Ú
make_ttgirÚtensordesc_meta)5r   rÒ   rÓ   rÔ   r   r×   Úadd_convert_to_ttgpuirr"   r.   rS   r2   rá   ÚttgpuirÚadd_coalesceÚadd_f32_dot_tcÚadd_remove_layout_conversionsÚadd_optimize_thread_localityr	   Úadd_accelerate_matmulrC   rD   Úadd_optimize_epilogueÚadd_optimize_dot_operandsÚadd_hoist_layout_conversionsÚadd_fuse_nested_loopsrÕ   rÚ   rÞ   r
   r#   r$   Úadd_schedule_loopsr1   Úadd_pipelineÚadd_coalesce_async_copyrL   Ú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è   Ú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°·\±\À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ÔOÞÜ�J‰J×Ñ×6Ñ6°r¿<¹<ÔHÜ�‰×'Ñ'¨Ô+Ø× Ñ ×&Ñ&Ó(¨FÓ2Ø×-Ñ-×3Ñ3°CÖ8�Ü—
‘
×"Ñ"×AÑAÀ"ÖKñ 9ä�‰×4Ñ4°RÔ8Ü�‰×2Ñ2°2Ô6Ü)¨'¯,©,×7Ñ7Ü�J‰J×Ñ×6Ñ6°rÔ:Ü�N‰N×8Ñ8¸Ô<Ü�
‰
×Ñ×3Ñ3°BÔ7Þ '×"4Ñ"4°qÓ"8Ü�J‰J×Ñ×1Ñ1°"×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                 ó(  • U n[         R                  " UR                  5      nUR                  5         [        R
                  R                  U5        [        R
                  R                  U5        [        R                  R                  U5        [        R                  R                  U5        [        R
                  R                  U5        [        R                  R                  U5        UR                  US5        UR!                  5       US'   U$ )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	  ÚHIPBackend.gluon_to_ttgir  s·   € àˆÜ�_Š_˜SŸ[™[Ó)ˆØ
�‰Ôä�‰× Ñ  Ô$Ü�‰×/Ñ/°Ô3Ü�‰×Ñ˜rÔ"Ü�‰×&Ñ& rÔ*Ü�‰×&Ñ& rÔ*Ü�‰×7Ñ7¸Ô;à
�‰ˆsÐ$Ô%Ø&)×&AÑ&AÓ&CˆÐ"Ñ#Øˆ
r   c                 óš  • U n[         R                  " UR                  5      nUR                  5         [        R
                  R                  R                  XBR                  5        Sn[        R
                  R                  R                  XBR                  U5        [
        R                  R                  U5        [
        R                  R                  U5        [
        R                  R                  U5        [        R
                  R                  R                  U5        [         R"                  (       a*  [         R"                  R%                  SXCR                  5        Sn[        R
                  R                  R'                  XBR                  U5        [
        R(                  R+                  U5        [
        R(                  R-                  U5        [
        R                  R/                  U5        [
        R                  R1                  U5        [
        R(                  R+                  U5        [
        R(                  R-                  U5        [
        R(                  R3                  U5        UR4                  R7                  5       S:w  a>  [        R
                  R                  R9                  XBR                  UR:                  5        [         R"                  (       a*  [         R"                  R%                  SXCR                  5        [<        R>                  R@                  (       d>  [<        R>                  RB                  (       d  [
        RD                  RG                  U5        [        R
                  R                  RI                  XF5        URK                  US5        [<        R>                  RB                  (       aá  [<        R>                  R@                  (       da  [         R                  " UR                  5      nUR                  5         [
        RD                  RG                  U5        URK                  US5        [         R                  " UR                  5      nUR                  5         [
        RD                  RM                  U5        URK                  US5        [N        RP                  " 5         [N        R                  " 5       n[N        RR                  " X75      n[        RT                  " U5        S	n	[<        R>                  RV                  (       a  S
n	[N        RX                  " U[        RZ                  UR                  U	5        [        R\                  " X‚R                  5        [        R^                  " US5        [        R`                  " USS5        [        R`                  " USS5        [        R`                  " USS5        [        R`                  " USURb                  S:H  5        URe                  5        V
s/ s H  oªRg                  5       (       a  M  U
PM     nn
US   Ri                  [        Rj                  5        US   Rm                  SSURn                  URb                  -   35        SUR4                  Rq                  S5      ;   a  US   Rm                  SS5        US   Rm                  SS5        US   Rm                  SURr                   SURr                   35        URt                  (       a  SOSnUS   Rm                  SU5        [<        R>                  RV                  (       a'  US   Rw                  S
5        US   Ry                  5         [        Rz                  " US   5        [<        R>                  RV                  (       a\  [}        [~        5      R€                  S-  n[ƒ        US -  5      [ƒ        US!-  5      [ƒ        US"-  5      /n[N        R„                  " XŽ5        OtUR†                  (       ac  UR†                   VVs/ s H%  u  nn[        Rˆ                  " X�5      (       d  M#  UPM'     nnn[‹        U5      S:”  a  [N        R„                  " XŽ5        [N        RŒ                  " U[N        RŽ                  UR                  S	/ UR�                  5        [        R’                  " UR                  5      (       a<  US   R•                  S#5        US   R•                  S$5        US   R•                  S%5        [<        R                  R–                  (       a  [        R˜                  " US   5        U R›                  S&5      US''   U R›                  S(5      =(       d    SUS)'   U R›                  S*5      =(       d    S+US,'   [        Rœ                  " U5        [        Rž                  " U5        [ƒ        U5      $ s  sn
f s  snnf )-Nr   Úttgpuir_to_llvmirTrK   Úllvmir_to_llvmÚ	make_llirzmake_llir.disable_line_infoz,make_llir.dump_ir_extract_di_local_variablesrI   ú+xnackiô  Ú__oclc_finite_only_optFÚ__oclc_correctly_rounded_sqrt32Ú__oclc_unsafe_math_optÚ__oclc_wavefrontsize64rR   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-f32rT   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_memoryr~   r·   ÚpatchÚadd_to_llvmirrÕ   rÚ   rÝ   Úadd_cf_to_llvmirÚadd_arith_to_llvmirrß   rL   rø   Úlower_instruction_sched_hintsr1   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_constantrS   Úget_functionsÚis_declarationÚset_calling_convÚCALLING_CONV_AMDGPU_KERNELÚadd_fn_attrr.   rù   r/   rE   Úadd_fn_target_featureÚadd_fn_asan_attrÚset_all_fn_arg_inregr   r\   r]   r_   Úlink_extern_libsr3   Úneed_extern_libÚlenÚoptimize_moduleÚOPTIMIZE_O3rA   Ú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_moderd   Úpathsrq   Úpaths                    r   r  ÚHIPBackend.make_llir  sÎ  € àˆä�_Š_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�
‰
×Ñ×6Ñ6°r¿<¹<ÔHð ˆÜ�
‰
×Ñ×1Ñ1°"·l±lÀOÔTÜ�‰×$Ñ$ RÔ(Ü�‰× Ñ  Ô$Ü�‰×*Ñ*¨2Ô.ä�
‰
×Ñ×5Ñ5°bÔ9ä×%×%Ü×&Ñ&×,Ñ,Ð-@À"ÇkÁkÔRð ˆ	Ü�
‰
×Ñ×(Ñ(¨¯\©\¸9ÔEÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!ä�‰×'Ñ'¨Ô+Ü�‰×*Ñ*¨2Ô.Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(à× Ñ ×&Ñ&Ó(¨FÓ2Ü�J‰J×Ñ×<Ñ<¸RÇÁÈw×OaÑOaÔbô ×%×%Ü×&Ñ&×,Ñ,Ð-=¸rÇ;Á;Ô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ô 	×ÒÔÜ—,’,“.ˆÜ—>’> #Ó/ˆÜ× Ò  Ô*ØˆÜ×Ñ×(×(Ø&ˆOÜ×Ò˜x¬×):Ñ):¸G¿L¹LÈ/ÔZô 	×Ò˜H§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Ò4�b×<MÑ<M×<O�rÑ4ˆÐ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ô
 ×!Ò! (Õ2Ø× × Ø.5×.AÒ.AÔiÒ.A™l˜t TÄS×EXÒEXÐYa×Eh—TÑ.AˆEÑiÜ�5‹z˜A‹~Ü×%Ò% hÔ6ä×Ò˜X¤t×'7Ñ'7¸¿¹ÀrÈ2Èw×OgÑOgÔhô ×$Ò$ W§\¡\×2Ñ2Ø�‰F×!Ñ!Ð"<Ô=Ø�‰F×!Ñ!Ð"<Ô=Ø�‰F×!Ñ!Ð"<Ô=ä�9‰9×*×*Ü×3Ò3°C¸±FÔ;ð !×-Ñ-¨lÓ;ˆ�ÑØ+.×+;Ñ+;Ð<]Ó+^×+cÐbcˆÐ'Ñ(Ø,/×,<Ñ,<Ð=cÓ,d×,iÐhiˆÐ(Ñ)ä×$Ò$ XÔ.ô 	× Ò  Ô*Ü�8‹}ÐùòA QùóJ js   ØgØgà"gà9gc           	      óæ  • [         R                  " SU 5      n[        U5      S:X  d   eUS   US'   / nSUR                  ;   a  SOSn[        R
                  " U R                  S5      5      R                  5       nUS   S	-   U-   n[        R                  " U [        R                  UR                  XTUR                  U5      n[        R                  " U [        R                  UR                  XTUR                  U5        [        R                  " U [        R                  UR                  XTUR                  S
5      n	[        R                  R                   (       a  [#        S5        [#        U	5        U	$ )Nz3define amdgpu_kernel void @([a-zA-Z_][a-zA-Z0-9_]*)r   r   rq   Úgfx11z-real-true16rI   rj   rh   Fz!// -----// AMDGCN Dump //----- //)ÚreÚfindallrB  r"   rm   rn   ro   rp   r   Útranslate_to_mirr	   r4  rA   Údump_sched_dagÚtranslate_to_asmr
   Údump_amdgcnÚprint)
r  r©   r�   ÚnamesÚflagsÚfeaturesÚir_hashÚdump_file_idrh   Úamdgcns
             r   Úmake_amdgcnÚHIPBackend.make_amdgcnº  s'  € ô
 —
’
ÐQÐSVÓWˆÜ�5‹z˜Q‹ÐˆØ  ™8ˆ�ÑàˆØ%,°·±Ó%<‘>À"ˆÜ—.’. §¡¨GÓ!4Ó5×?Ñ?ÓAˆØ˜Q‘x #‘~¨Ñ/ˆÜ×!Ò! #¤s×'8Ñ'8¸'¿,¹,ÈÐY`×YqÑYqØ".ó0ˆä×Ò˜C¤×!2Ñ!2°G·L±LÀ(ÐSZ×SkÑSkØ(ô	*ä×&Ò& s¬C×,=Ñ,=¸w¿|¹|ÈXÐ^e×^vÑ^vØ',ó.ˆä�9‰9× × ÜÐ5Ô6Ü�&ŒMØˆr   c                 ó¤  • Sn[         R                  R                  (       a  Sn[        R                  " XR
                  U5      n[        R                  " 5        n[        R                  " 5        n[        UR                  S5       nUR                  U5        S S S 5        [        R                  " UR                  UR                  5        S S S 5        [        UR                  S5       nUR                  5       n	S S S 5        S S S 5        W	$ ! , (       d  f       Nz= f! , (       d  f       NX= f! , (       d  f       N:= f! , (       d  f       W	$ = f)NrI   r  ÚwbÚrb)r
   r(  r2  r	   Úassemble_amdgcnr"   ÚtempfileÚNamedTemporaryFileÚopenrq   ÚwriteÚ
link_hsacoÚread)
r  r©   r�   rO  r†   Útmp_outÚtmp_inÚfd_inÚfd_outrÉ   s
             r   Ú
make_hsacoÚ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   ÁE Á0DÂDÂ3DÃE Ã*D/Ã;E Ä
DÄDÄ
D,	Ä(E Ä/
D=	Ä9E Å 
Ec                 óP  ^ ^• U[         R                  :X  a  UU 4S jUS'   UU 4S jUS'   OU[         R                  :X  a
  UU 4S jUS'   UU 4S jUS'   UU 4S jUS	'   UU 4S
 jUS'   [        R                  R
                  b$  [        R                  R                  T UTUS 5        g g )Nc                 ó(   >• TR                  XT5      $ r§   )rÑ   ©r  r©   r�   rb   s     €€r   r   Ú'HIPBackend.add_stages.<locals>.<lambda>ã  s   ø€ °4·>±>À#ÐQXÔ3Yr   r×   c                 ó(   >• TR                  XT5      $ r§   )rè   ry  s     €€r   r   rz  ä  s   ø€ °D·O±OÀCÐSZÔ4[r   Úttgirc                 ó(   >• TR                  XT5      $ r§   )r	  ry  s     €€r   r   rz  æ  s   ø€ °D×4GÑ4GÈÐW^Ô4_r   c                 ó(   >• TR                  XT5      $ r§   )r  ry  s     €€r   r   rz  ç  s   ø€ ¨t¯~©~¸cÈWÔ/Ur   Úllirc                 ó(   >• TR                  XT5      $ r§   )re  ry  s     €€r   r   rz  è  s   ø€ °×1AÑ1AÀ#ÐQXÔ1Yr   rd  c                 ó(   >• TR                  XT5      $ r§   )ru  ry  s     €€r   r   rz  é  s   ø€ °·±ÀÈwÔ0Wr   r†   )r   ÚTRITONÚGLUONr
   r•   Úadd_stages_inspection_hook)rb   Ústagesr�   r�   s   ` ` r   Ú
add_stagesÚ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                 ó   • U R                    $ r§   r   )rb   s    r   rt   ÚHIPBackend.hashí  s   € à—+‘+�Ðr   )rŠ   )#rv   rw   rx   ry   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  re  ru  r†  Ú	functoolsÚ	lru_cachert   r|   Ú__classcell__)r‹   s   @r   r~   r~   f   sQ  ø† Ø€OØ,1Ð)àð' 	ó 'ó ð'ð"˜yð "¨T÷ "ð
%¨#ô %ð" Sô "ò2
ò?ð>  S¨* _Ñ 5ô >ò
:ð
 ñó ðð ñó ðð ñó ðð ñó ðð  ñ<ó ð<ð| ñó ðð  ñZó ðZðx ñó ðð. ñó ðò
\ð ×ÒÓñ ó ö r   r~   )Útriton.backends.compilerr   r   r   Útriton._C.libtritonr   r   r   r	   Útritonr
   Údataclassesr   Útypingr   r   r   Útypesr   rm   rk  rX  rŒ  rZ   Úpathlibr   r   r$   r(   r+   r~   r   r   r   Ú<module>r–     sz   ðß EÑ Eß 5Ó 5Ý Ý !ß #Ñ #Ý Û Û Û 	Û Û Ý ð0˜Yô 0òXò
rñ �$Ñ÷D?ð D?ó ðD?ôNI �õ I r   