+
    Š°³j¡j  ã                   ó0  € ^ RI HtHtHt ^ RIHtHtHtHt ^ RI	H
t
 ^ RIHt ^ RIHtHtHt ^ RIHt ^ RIt^ RIt^ RIt^ RIt^ RIt^ RIt^ RIHt R	 R
 ltR tR tR tR tR t R t!R t"R t#R t$R t%]! RR7       ! R R4      4       t& ! R R]4      t'R# )é    )ÚBaseBackendÚ	GPUTargetÚLanguage)ÚirÚpassesÚllvmÚamd)Úknobs)Ú	dataclass)ÚAnyÚDictÚTuple)Ú
ModuleTypeN)ÚPathc                ó$   € V ^8„  d   QhR\         /# ©é   Útarget©r   )Úformats   "Ú]/home/luca/screen-checking/.venv/lib/python3.14/site-packages/triton/backends/amd/compiler.pyÚ__annotate__r      s   € ÷ 0ñ 0œYñ 0ó    c                 ó   € R  # )c                 ó   € R# )é   )r   r   r   © )Úlhs_typeÚrhs_types   &&r   Ú<lambda>Ú"get_min_dot_size.<locals>.<lambda>   s   € ¡ir   r   ©r   s   &r   Úget_min_dot_sizer#      s
   € ñ 0Ð/r   c                 óª   € \         P                  P                  f   V R8H  ;'       g    V R8H  ;'       d    VRJ # \         P                  P                  # )NÚgfx942Úgfx950T)r
   r	   Úuse_block_pingpong)ÚarchÚuse_async_copys   &&r   Úis_pingpong_schedule_enabledr*      sO   € ä�9‰9×'Ñ'Ò/ð �HÑ×MÐM ¨Ñ!1×!LÐ!L°nÈÐ6Lð RÜ5:·Y±Y×5QÑ5QðRr   c                 óÆ   € \         P                  P                  f-   V R8H  ;'       g     RV 9   ;'       g    RV 9   ;'       g    RV 9   # \         P                  P                  # )Nr%   Úgfx110Úgfx115Úgfx120)r
   r	   Úuse_in_thread_transpose©r(   s   &r   Úis_in_thread_transpose_enabledr1      s]   € ä�9‰9×,Ñ,Ò4ð �HÑ×XÐX ¨DÑ 0×XÐX°HÀÑ4D×XÐXÈÐTXÑHXð \Ü:?¿)¹)×:[Ñ:[ð\r   c                 óx   € \         P                  P                  f   V R9   # \         P                  P                  # )N)r&   Úgfx1250)r
   r	   r)   r0   s   &r   Úis_async_copy_enabledr4       s/   € Ü.3¯i©i×.FÑ.FÒ.NˆDÐ)Ñ)ÐlÔTY×T]ÑT]×TlÑTlÐlr   c                 ó   € V R9   # ©r3   ©r3   r   r0   s   &r   Úis_coexec_scheduler_supportedr8   $   ó   € Ø�;ÑÐr   c                 óx   € \         P                  P                  e   \         P                  P                  # V R9   # ©Nr7   )r
   r	   Úuse_coexec_schedulerr0   s   &r   Úis_coexec_scheduler_enabledr=   (   s-   € Ü‡y�y×%Ñ%Ò1Ü�y‰y×-Ñ-Ð-Ø�;ÑÐr   c                 óx   € \         P                  P                  e   \         P                  P                  # V R9   # r;   )r
   r	   Úuse_expert_schedulingr0   s   &r   Úis_expert_scheduling_enabledr@   .   s-   € Ü‡y�y×&Ñ&Ò2Ü�y‰y×.Ñ.Ð.Ø�;ÑÐr   c                 ó   € V R9   # )r%   )r%   r&   r3   r   r0   s   &r   Úis_fpsan_supportedrB   4   s   € ØÐ2Ñ2Ð2r   c                 ó   € V R9   # r6   r   r0   s   &r   Úis_consan_supportedrD   8   r9   r   c                 ó8   € V P                  R 4      '       d   R# R# )Úgfx11z-real-true16Ú )Ú
startswithr0   s   &r   Údisable_real_true16_featurerI   <   s   € Ø!Ÿ_™_¨W×5Ò5ˆ>Ð=¸2Ð=r   c                 ó>  € \        V \        4      '       g   \        V 4      # . pV P                  R 4       F[  pVP	                  R4      w  r4pVP                  4       pV'       g   K1  TP                  Y4'       d   VP                  4       MR34       K]  	  \        V4      # )Ú,Ú=rG   )Ú
isinstanceÚstrÚtupleÚsplitÚ	partitionÚstripÚappend)ÚattrsÚparsedÚattrÚnameÚsepÚvalues   &     r   Ú_parse_llvm_fn_attrsrZ   @   s{   € Ü�eœS×!Ò!Ü�U‹|ÐØ€FØ—‘˜CÖ ˆØŸ>™>¨#Ó.Ñˆ�5Ø�z‰z‹|ˆß‰4Ø�M‰M˜4¶# §¡¤¸2Ð>Ö?ñ	 !ô
 �‹=Ðr   T)Úfrozenc                   ó�   a € ] tR t^Lt o ^t^ t^t^tRtRt	Rt
RtRtR
tRtRtRtRt^ t^tRt^ tRtRtRtRtR tR tV 3R	 ltR
tV tR# )Ú
HIPOptionsNFTÚieeeÚhiprG   c                ó.  € \        V P                  ^R 4      pV^
8¼  d   ^ M^@p\        P                  V RV4       V P                  ^ 8”  d*   V P                  V P                  ^,
          ,          ^ 8X  g   Q R4       hV P                  R8X  dM   V P
                  ^8w  d<   \        P                  ! RV P
                   R24       \        P                  V R^4       \        P                  V R\        V P                  4      4       \        \        4      P                  R,          pV P                  f   / M\        V P                  4      pR F  p\        W5 R
2,          4      WE&   K  	  \        P                  V R\!        VP#                  4       4      4       R	# )é   Ú	warp_sizez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.ÚkpackÚllvm_fn_attrsÚlibNz.bcÚextern_libséþÿÿÿ)ÚocmlÚockl)Úintr(   ÚobjectÚ__setattr__Ú	num_warpsrc   ÚwarningsÚwarnrZ   rd   r   Ú__file__Úparentrf   ÚdictrN   rO   Úitems)ÚselfÚ	gfx_majorrb   Údefault_libdirrf   re   s   &     r   Ú__post_init__ÚHIPOptions.__post_init__p   sZ  € Ü˜Ÿ	™	 ! 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ä×Ñ˜4 Ô2FÀt×GYÑGYÓ2ZÔ[äœh›×.Ñ.°Õ6ˆØ ×,Ñ,Ò4‘b¼$¸t×?OÑ?OÓ:PˆÛ#ˆCÜ" >°e¸3°KÕ#?Ó@ˆKÓñ $ä×Ñ˜4 ´°k×6GÑ6GÓ6IÓ0JÖKr   c           	     óø   € R P                  V P                  P                  4        UUu. uF  w  rV RV 2NK  	  upp4      p\        P                  ! VP                  R4      4      P                  4       # u uppi )Ú_Ú-úutf-8)ÚjoinÚ__dict__rs   ÚhashlibÚsha256ÚencodeÚ	hexdigest)rt   rW   ÚvalÚkeys   &   r   ÚhashÚHIPOptions.hash…   sa   € Ø�h‰h¸¿¹×9LÑ9LÔ9NÔOÑ9N©I¨D˜4˜&  # ›Ñ9NÒOÓPˆÜ�~Š~˜cŸj™j¨Ó1Ó2×<Ñ<Ó>Ð>ùó Ps   ©A6
c                óŒ  <€ V ^8„  d   Qh/ S[ ;R&   S[ ;R&   S[ ;R&   S[ ;R&   S[;R&   S[;R&   S[;R&   S[;R&   S[S[,          ;R	&   S[S[,          ;R
&   S[;R&   S[S[,          ;R&   S[;R&   S[;R&   S[ ;R&   S[ ;R&   S[;R&   S[ ;R&   S[;R&   S[;R&   S[;R&   S[S[S[S[S[3,          R3,          ,          ;R&   # )r   rm   Úwaves_per_euÚ
num_stagesÚnum_ctasrf   ÚdebugÚsanitize_overflowr(   Úsupported_fp8_dtypesÚ!deprecated_fp8_dot_operand_dtypesÚdefault_dot_input_precisionÚallowed_dot_input_precisionsÚenable_fp_fusionÚlaunch_cooperative_gridÚmatrix_instr_nonkdimrc   Úallow_flush_denormÚmax_num_imprecise_acc_defaultÚbackend_nameÚinstrumentation_modeÚschedule_hint.rd   )rj   rr   ÚboolrN   r   )r   Ú__classdict__s   "€r   r   ÚHIPOptions.__annotate__L   sE  ø‡ ‚ áÑñ ñ Ññ ñ Ññ	 ñ
 Ññ ñ Ññ ñ Ññ ñ Ñ"ñ ñ Ññ ñ  ¡�*ÑRñ ñ (-©S¥zÑ6ñ ñ  "%Ñ-ñ! ñ" #(©¥*ÑKñ# ñ$ Ñ!ñ% ñ& "Ñ)ñ' ñ( Ñ!ñ) ñ* �Nñ+ ñ, Ñ$ñ- ñ. $'Ñ*ñ/ ñ0 Ññ1 ñ2 Ñ"ñ3 ñ: Ññ; ñD ™™u¡S©# X�°Ð3Õ4Õ4Ñ9òE r   r   )Úfp8e4nvÚfp8e5Úfp8e5b16Úfp8e4b8)r^   Úbf16x3Úbf16x6)Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__rm   rˆ   r‰   rŠ   rf   r‹   rŒ   r(   r�   rŽ   r�   r�   r‘   r’   r“   rc   r”   r•   r–   r—   r˜   rd   rw   r…   Ú__annotate_func__Ú__static_attributes__Ú__classdictcell__)rš   s   @r   r]   r]   L   s˜   ø‡ € à€IØ€LØ€JØ€HØ€KØ€EØ"ÐØ€Dð
 (SÐØ46Ð%Ø'-ÐØ/KÐ Ø!ÐØ$)ÐØ !ÐØ€EØ$ÐØ)*Ð!Ø€LØ "Ðð €Mð
 8:€MòLò*?÷s ƒ r   r]   c                   ó~  a a€ ] tR t^Št oRtRt]V3R lR l4       tV3R lV 3R lltV3R lR lt	V3R	 lR
 lt
R tR tV3R lR ltR tRt]R 4       t]R 4       t]R 4       t]R 4       t]R 4       t]R 4       t]R 4       t]R 4       t]R 4       tR t]P6                  ! 4       R 4       tV3R ltRtVtV ;t # )Ú
HIPBackendNFc                ó    <€ V ^8„  d   QhRS[ /# r   r   )r   rš   s   "€r   r   ÚHIPBackend.__annotate__�   s   ø€ ÷ 'ñ '¡	ñ 'r   c                ó    € V P                   R 8H  # )r_   )Úbackendr"   s   &r   Úsupports_targetÚHIPBackend.supports_targetŽ   s   € à�~‰~ Ñ&Ð&r   c                ó$   <€ V ^8„  d   QhRS[ RR/# )r   r   ÚreturnNr   )r   rš   s   "€r   r   r¬   ’   s   ø€ ÷ "ñ "™yð "¨Tñ "r   c                óx   <€ \         SV `  V4       \        VP                  \        4      '       g   Q hR V n        R# )ÚhsacoN)ÚsuperÚ__init__rM   r(   rN   Ú
binary_ext)rt   r   Ú	__class__s   &&€r   r¶   ÚHIPBackend.__init__’   s.   ø€ Ü‰Ñ˜Ô Ü˜&Ÿ+™+¤s×+Ò+Ð+Ð+Ø!ˆŽr   c                ó    <€ V ^8„  d   QhRS[ /# ©r   r²   )rN   )r   rš   s   "€r   r   r¬   —   s   ø€ ÷ %ñ %©#ñ %r   c                ó    € R VP                    2# )úhip:r0   ©rt   Úoptionss   &&r   Úget_target_nameÚHIPBackend.get_target_name—   s   € Ø�g—l‘l�^Ð$Ð$r   c                ó    <€ V ^8„  d   QhRS[ /# r»   )r   )r   rš   s   "€r   r   r¬   š   s   ø€ ÷ "ñ "¡Sñ "r   c                óÈ  a€ \         ;QJ d    V3R  lR 4       F  '       g   K   RM	  RM! V3R  lR 4       4      '       d   RSR&   RSR&   R\        P                  P                  ;'       g    V P                  P
                  /pSP                  R^4      ^8”  dS   \        P                  ! V P                  P
                  4      '       g#   \        RV P                  P
                   24      hV P                  P
                  R8X  dC   \        \        P                  4      pVP                  R	04       \        \        V4      4      VR
&   RS9  d&   \        \        \        P                   4      4      VR&   V P                  P
                  R8X  dD   \        \        P"                  4      pVP                  RR04       \        \        V4      4      VR&   RS9  d   \        P$                  P&                  VR&   TP                  \        P(                  P+                  4        Uu/ uF"  qUS9   g   K  SV,          f   K  VSV,          bK$  	  up4       \        R/ VB # u upi )c              3   óL   <"  € T F  qSP                  R R4      9   x € K  	  R# 5i)r—   rG   N)Úget)Ú.0ÚmodeÚoptss   & €r   Ú	<genexpr>Ú+HIPBackend.parse_options.<locals>.<genexpr>œ   s    øé € ÐSÉ
À�t—x‘xÐ 6¸Ó;Ö;Ë
ùs   ƒ!$TFr‹   rŒ   r(   rŠ   znum_ctas > 1 not supported on r%   Útf32r�   r�   r&   rž   rŸ   rŽ   r‘   )Úconsanr   )Úanyr
   ÚruntimeÚoverride_archr   r(   rÅ   r	   Úsupports_multi_cta_launchÚ
ValueErrorÚsetr]   r�   ÚupdaterO   Úsortedr�   rŽ   ÚlanguageÚdefault_fp_fusionÚ__dataclass_fields__Úkeys)rt   rÈ   Úargsr�   rŽ   Úks   &f    r   Úparse_optionsÚHIPBackend.parse_optionsš   sÛ  ø€ ç‹3ÔSÉ
ÓS�3�3Š3ÔSÉ
ÓS×SÒSØ ˆD�‰MØ(-ˆDÐ$Ñ%àœŸ™×3Ñ3×GÐ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ÐfgÕah”Z�Q˜˜Q�’ZÑ)OÑuÔvÜÑ!˜DÑ!Ð!ùò vs   È%IÈ2IÈ?Ic                óH   € VP                   VP                  VP                  3# ©N)rm   rŠ   Úshared)rt   Úmetadatas   &&r   Úpack_metadataÚHIPBackend.pack_metadata¸   s%   € à×ÑØ×ÑØ�O‰Oð
ð 	
r   c                ó0   € R \        V P                  4      /# )Úmin_dot_size)r#   r   r¾   s   &&r   Úget_codegen_implementationÚ%HIPBackend.get_codegen_implementation¿   s   € ØÔ 0°·±Ó =Ð>Ð>r   c                ó6   <€ V ^8„  d   QhRS[ S[S[3,          /# r»   )r   rN   r   )r   rš   s   "€r   r   r¬   Â   s   ø€ ÷ >ñ >¡¡S©* _Õ 5ñ >r   c                ó   € ^ RI Hp RV/# )r   )Ú	libdeviceztriton.language.extra.libdevice)Útriton.language.extra.hipré   )rt   ré   s   & r   Úget_module_mapÚHIPBackend.get_module_mapÂ   s   € Ý7à1°9Ð=Ð=r   c                ó    € \         P                  ! V4       \        P                  '       d"   \        P                  P                  V4       R # R # rÞ   )r	   Úload_dialectsrª   Úinstrumentation)rt   Úctxs   &&r   rî   ÚHIPBackend.load_dialectsÇ   s5   € Ü×Ò˜#ÔÜ×%×%Ð%Ü×&Ñ&×4Ñ4°SÖ9ñ &r   c                óÐ  € \         P                  f    ^ R IpR\         n        M\         P                  '       d   ^ R IpRp\	        V R4      '       d   V P                  4       V8*  # \         P                  '       dP   \        V XP                  4      '       d4   \	        V R4      '       d"   V P                  4       P                  4       V8*  # R#   \         d    R\         n         L¦i ; i)NTFÚ	ptr_rangeÚuntyped_storageiÿÿÿ)
rª   Ú_torch_availableÚtorchÚImportErrorÚhasattrró   rM   ÚTensorrô   Úsize)Úargrö   Ú
MAX_INT_32s   &  r   Úis_within_2gbÚHIPBackend.is_within_2gbÒ   sµ   € ä×&Ñ&Ò.ð4ÛØ.2”
Õ+ô ×(×(Ð(Ûàˆ
Ü�3˜×$Ò$Ø—=‘=“? jÑ0Ð0Ü×&×&Ð&¬:°c¸5¿<¹<×+HÒ+HÌWÐUXÐZk×MlÒMlØ×&Ñ&Ó(×-Ñ-Ó/°:Ñ=Ð=Ùøô ô 4Ø.3”
Ö+ð4ús   ”C ÃC%Ã$C%c                óX   € \         P                  ! V 4      pR V 9   d   VR^ ..,          pV# )ÚSztt.pointer_range)r   Ú
parse_attr)ÚdescÚrets   & r   r  ÚHIPBackend.parse_atträ   s1   € ä×$Ò$ TÓ*ˆØ�$Œ;ØÐ'¨Ð,Ð-Õ-ˆCØˆ
r   c                ó¼   € \         P                  ! V 3/ VB p\        P                  P                  '       d%   \
        P                  V 4      '       d
   VR ,          pV# )r   )r   Úget_tensor_specializationr
   r	   Úuse_buffer_opsrª   rý   )rû   Úkwargsr  s   &, r   r  Ú$HIPBackend.get_tensor_specializationë   sE   € ä×3Ò3°CÑB¸6ÑBˆÜ�9‰9×#×#Ð#¬
×(@Ñ(@À×(EÒ(EØ�3�JˆCØˆ
r   c                ó  € \         P                  ! V P                  4      pVP                  4        \        P
                  P                  V4       \        P                  ! VP                  4      '       g    \        P                  P                  V4       \        P
                  P                  V4       \        P                  P                  V4       \        P                  P                  V4       \        P
                  P                  V4       \        P                  P!                  V4       \        P
                  P#                  V4       \        P                  P%                  V4       VP'                  V R 4       V # ©Ú	make_ttir)r   Úpass_managerÚcontextÚenable_debugr   ÚcommonÚadd_inlinerr	   Úsupports_tdmr(   ÚttirÚ(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Ÿ[™[Ó)ˆØ
�‰ÔÜ�‰×!Ñ! "Ô%Ü×Ò §¡×-Ò-Ü�K‰K×@Ñ@ÀÔDÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ Ô#Ü�‰×)Ñ)¨"Ô-Ü�‰×Ñ˜bÔ!Ü�‰×#Ñ# BÔ'Ü�‰×$Ñ$ RÔ(Ü�‰×#Ñ# BÔ'Ø
�‰ˆs�KÔ Øˆ
r   c                ó¸  € \         P                  ! V P                  4      pVP                  4        \        P
                  P                  VR VP                   2VP                  VP                  VP                  4       VP                  V R4       \         P                  ! V P                  4      pVP                  4        Rp\        P                  P                  V4       \        P                  P                  W44       \        P                  P                  V4       \        P                  P!                  V4       \"        P                  P                  P%                  W2P                  VP&                  VP(                  4       \        P                  P                  V4       \"        P                  P                  P+                  V4       \"        P                  P                  P-                  W2P                  4       \"        P                  P                  P/                  V4       \"        P                  P                  P1                  V4       \        P                  P3                  V4       \        P4                  P7                  V4       \        P
                  P9                  V4       \        P4                  P7                  V4       \;        VP                  4      p\=        VP                  V4      p\"        P                  P                  P?                  V4       \"        P                  P                  PA                  W2PB                  4       \"        P                  P                  PE                  W5V4       V'       d4   \"        P                  P                  PG                  W2P                  4       \"        P                  P                  PI                  V4       \        P4                  P7                  V4       \        P                  P                  V4       \        P                  PK                  V4       \M        VP                  4      '       dI   \"        P                  P                  PO                  V4       \        P                  P                  V4       \"        P                  P                  PQ                  V4       V'       dE   VPB                  ^8”  d4   \"        P                  P                  PS                  W2PB                  4       \T        P"                  PV                  '       dØ   \"        P                  P                  PY                  V4       \        P4                  P7                  V4       \"        P                  P                  P[                  VVP                  \T        P"                  P\                  \T        P"                  P^                  4       \"        P                  P                  Pa                  V4       \"        P                  P                  Pc                  V4       \"        P                  P                  Pe                  V4       \        P4                  P7                  V4       \        P4                  Pg                  V4       \        P4                  Pi                  V4       VPj                  R8X  dd   \m        VP                  4      '       dI   \"        P                  P                  Po                  V4       \        P                  Po                  V4       VP                  V R4       V Pq                  4       VR&   V # )r½   Úmake_ttgir_earlyFÚfpsanÚ
make_ttgirÚtensordesc_meta)9r   r  r  r  r   r  Úadd_convert_to_ttgpuirr(   rm   rb   rŠ   r  ÚttgpuirÚadd_coalesceÚadd_f32_dot_tcÚadd_remove_layout_conversionsÚadd_optimize_thread_localityr	   Úadd_accelerate_matmulr“   rc   Úadd_optimize_epilogueÚadd_optimize_dot_operandsÚadd_hoist_layout_conversionsÚadd_sink_layout_conversionsÚadd_fuse_nested_loopsr  r  r  r4   r*   Ú add_optimize_descriptor_encodingÚadd_schedule_loopsr‰   Úadd_pipelineÚadd_coalesce_async_copyÚadd_convert_to_tensor_opsÚadd_reduce_data_duplicationr1   Úadd_in_thread_transposeÚadd_move_up_prologue_loadsÚadd_block_pingpongr
   r  Úadd_canonicalize_pointersÚadd_convert_to_buffer_opsÚuse_buffer_atomicsÚ%buffer_ops_analyze_small_tensor_rangeÚadd_optimize_buffer_op_ptrÚadd_fold_true_cmpiÚadd_prepare_if_combiningr  r  r—   rB   Úadd_fp_sanitizerÚget_tensordesc_metadata)r  rà   r¿   r  ÚemuTF32r)   r'   s   &&&    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¸Ô;Ü�
‰
×Ñ×6Ñ6°rÔ:ä�‰×,Ñ,¨RÔ0Ü�‰×'Ñ'¨Ô+Ü�‰×#Ñ# BÔ'Ü�‰×'Ñ'¨Ô+ä.¨w¯|©|Ó<ˆÜ9¸'¿,¹,ÈÓWÐÜ�
‰
×Ñ×;Ñ;¸BÔ?Ü�
‰
×Ñ×-Ñ-¨b×2DÑ2DÔEÜ�
‰
×Ñ×'Ñ'¨Ð<NÔOßÜ�J‰J×Ñ×6Ñ6°r¿<¹<ÔHÜ�
‰
×Ñ×4Ñ4°RÔ8Ü�‰×'Ñ'¨Ô+Ü�‰×4Ñ4°RÔ8Ü�‰×2Ñ2°2Ô6Ü)¨'¯,©,×7Ò7Ü�J‰J×Ñ×6Ñ6°rÔ:Ü�N‰N×8Ñ8¸Ô<Ü�
‰
×Ñ×5Ñ5°bÔ9ß '×"4Ñ"4°qÔ"8Ü�J‰J×Ñ×1Ñ1°"×6HÑ6HÔIä�9‰9×#×#Ð#Ü�J‰J×Ñ×8Ñ8¸Ô<Ü�M‰M×+Ñ+¨BÔ/Ü�J‰J×Ñ×8Ñ8ØØ—‘Ü—	‘	×,Ñ,Ü—	‘	×?Ñ?ô	ô �J‰J×Ñ×9Ñ9¸"Ô=ä�
‰
×Ñ×-Ñ-¨bÔ1Ü�
‰
×Ñ×3Ñ3°BÔ7Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(Ø×'Ñ'¨7Ô2Ô7IÈ'Ï,É,×7WÒ7WÜ�J‰J×Ñ×/Ñ/°Ô3Ü�N‰N×+Ñ+¨BÔ/Ø
�‰ˆs�LÔ!Ø&)×&AÑ&AÓ&CˆÐ"Ñ#Øˆ
r   c                óÞ  € T p\         P                  ! VP                  4      pVP                  4        \        P
                  P                  V4       \        P
                  P                  V4       \        P                  P                  V4       \        P                  P                  V4       \        P
                  P                  V4       \        P                  P                  V4       \        P                  P                  V4       \         P                  P                  P#                  V4       \        P                  P%                  V4       VP&                  R 8X  dd   \)        VP*                  4      '       dI   \         P                  P                  P-                  V4       \        P                  P-                  V4       VP/                  VR4       VP1                  4       VR&   V# )r"  Úgluon_to_ttgirr$  )r   r  r  r  r   Úgluonr  Úadd_resolve_auto_encodingsr  Úadd_sccpr  Úadd_loop_aware_cser  r  r&  Ú add_combine_tensor_select_and_ifr	   Úadd_warp_pipelineÚadd_allocate_warp_groupsr—   rB   r(   rA  r  rB  )Úsrcrà   r¿   r  r  s   &&&  r   rF  ÚHIPBackend.gluon_to_ttgirF  s:  € àˆÜ�_Š_˜SŸ[™[Ó)ˆØ
�‰Ôä�‰× Ñ  Ô$Ü�‰×/Ñ/°Ô3Ü�‰×Ñ˜rÔ"Ü�‰×&Ñ& rÔ*Ü�‰×&Ñ& rÔ*Ü�‰×#Ñ# BÔ'Ü�‰×7Ñ7¸Ô;Ü�
‰
×Ñ×,Ñ,¨RÔ0Ü�‰×/Ñ/°Ô3à×'Ñ'¨7Ô2Ô7IÈ'Ï,É,×7WÒ7WÜ�J‰J×Ñ×/Ñ/°Ô3Ü�N‰N×+Ñ+¨BÔ/à
�‰ˆsÐ$Ô%Ø&)×&AÑ&AÓ&CˆÐ"Ñ#Øˆ
r   c           
     ó4  € T p\         P                  ! VP                  4      pVP                  4        \        P
                  P                  P                  WBP                  4       \        P
                  P                  P                  WBP                  4       \
        P                  P                  V4       \
        P                  P                  V4       \
        P                  P                  V4       R VP                  9   d<   \!        VP                  4      '       d!   \
        P                  P#                  VR4       \        P
                  P                  P%                  WBP                  4       R VP                  9   dy   \!        VP                  4      '       d^   \
        P                  P'                  V4       \
        P                  P)                  V4       \
        P*                  P-                  V4       \
        P                  P/                  V4       \0        P2                  '       d+   \0        P2                  P5                  RWCP                  4       \
        P                  P/                  V4       Rp\        P
                  P                  P7                  WBP                  V4       \        P
                  P                  P9                  WBP                  4       \
        P*                  P)                  V4       \
        P*                  P-                  V4       \
        P                  P;                  V4       \
        P                  P=                  V4       \
        P*                  P)                  V4       \
        P*                  P-                  V4       \
        P*                  P?                  V4       \0        P2                  '       d+   \0        P2                  P5                  RWCP                  4       \@        PB                  PD                  '       g@   \@        PB                  PF                  '       g    \
        PH                  PK                  V4       \        P
                  P                  PM                  WBP                  V4       \
        P                  PO                  V4       VPQ                  VR4       \@        PB                  PF                  '       dã   \@        PB                  PD                  '       gb   \         P                  ! VP                  4      pVP                  4        \
        PH                  PK                  V4       VPQ                  VR4       \         P                  ! VP                  4      pVP                  4        \
        PH                  PS                  V4       VPQ                  VR4       \T        PV                  ! 4        \T        P                  ! 4       p\T        PX                  ! W64      p\        PZ                  ! V4       Rp\@        PB                  P\                  '       d   R	p\T        P^                  ! V\        P`                  VP                  V4       \        Pb                  ! WrP                  4       \        Pd                  ! VR
4       \        Pf                  ! VRR4       \        Pf                  ! VRR4       \        Pf                  ! VRR4       \        Pf                  ! VRVPh                  ^@8H  4       VPk                  4        U	u. uF  q™Pm                  4       '       d   K  V	NK  	  p
p	\o        R V
 4       R4      pV'       g   \q        R4      hVPs                  \        Pt                  4       VR,          pVPw                  RV R24       VPx                  pV P{                  R4      pVe   TpVPw                  RRWÒPh                  ,           24       VPw                  RR4       VP|                  ^ 8w  d,   VPw                  RVP|                   RVP|                   24       \        VP                  4      '       d$   VPx                  ^8:  d   VPw                  RR4       VP€                  '       d   RMR pVPw                  R!V4       \@        PB                  P\                  '       d"   VPƒ                  R	4       VP…                  4        VP†                   F)  w  ppVP‰                  V4       VPw                  VV4       K+  	  VP                  R"8w  dW   \@        P                  PŠ                  '       g7   \@        P                  PŒ                  '       g   \        PŽ                  ! V4       \@        PB                  P\                  '       dn   \‘        \’        4      P”                  R#,          p\—        VR$,          4      \—        VR%,          4      \—        VR&,          4      .p\T        P˜                  ! VV4       MxVPš                  '       dg   VPš                   UUu. uF'  w  pp\        Pœ                  ! VV4      '       g   K%  VNK)  	  ppp\Ÿ        V4      ^ 8”  d   \T        P˜                  ! VV4       \T        P                   ! V\T        P¢                  VP                  R. VP¤                  RR'7       \        P¦                  ! VP                  4      '       d4   VP‰                  R(4       VP‰                  R)4       VP‰                  R*4       \@        P                  P¨                  '       d   \        Pª                  ! V4       WÑR+&   V P{                  R,4      VR-&   V P{                  R.4      VR/&   V P{                  R04      VR1&   V P{                  R24      ;'       g    ^ VR3&   V P{                  R44      ;'       g    ^VR5&   \        P¬                  ! V4       \        P®                  ! V4       \—        V4      # u up	i u uppi )6rÌ   r	   Úttgpuir_to_llvmirTÚ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_wavefrontsize64c              3   óT   "  € T F  qP                  4       '       g   K  Vx € K   	  R # 5irÞ   )Úis_external_linkage)rÆ   Úfns   & r   rÉ   Ú'HIPBackend.make_llir.<locals>.<genexpr>À  s   é € ÐG¡s ×.DÑ.D×.FŸ"š"£sùs   ‚(ž
(NzCould not find kernel functionrŠ   zamdgpu-cluster-dimsz,1,1zttg.total-num-warpszamdgpu-flat-work-group-sizez1,zuniform-work-group-sizeÚtruezamdgpu-waves-per-eurK   zamdgpu-sched-strategyÚcoexeczpreserve-signr^   zdenormal-fp-math-f32r3   re   z
asanrtl.bczocml.bczockl.bc)Údisable_vector_combinezamdgpu-no-workgroup-id-xzamdgpu-no-workgroup-id-yzamdgpu-no-workgroup-id-zrm   z
ttg.sharedrß   zttg.global_scratch_memory_sizeÚglobal_scratch_sizez#ttg.global_scratch_memory_alignmentÚglobal_scratch_alignzttg.profile_scratch_memory_sizeÚprofile_scratch_sizez$ttg.profile_scratch_memory_alignmentÚprofile_scratch_align)Xr   r  r  r  r	   r   r&  Úadd_update_async_wait_countr(   Úadd_warp_pipeline_conversionÚconvertÚadd_scf_to_cfrG  r  Úadd_index_to_llvmirr—   rD   Úadd_prepare_consan_capturesÚadd_allocate_shared_memoryÚadd_concurrency_sanitizerr  r  r  Ú"add_allocate_global_scratch_memoryrª   rï   ÚpatchÚadd_to_llvmirÚadd_warp_specialize_to_llvmÚadd_cf_to_llvmirÚadd_arith_to_llvmirr  r
   ÚcompilationÚdisable_line_infoÚ"dump_ir_extract_di_local_variablesÚllvmirÚadd_di_scopeÚadd_builtin_func_to_llvmirÚadd_reconcile_unrealized_castsr  Ú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_constantrb   Úget_functionsÚis_declarationÚnextÚRuntimeErrorÚset_calling_convÚCALLING_CONV_AMDGPU_KERNELÚadd_fn_attrrm   Úget_int_attrrˆ   r=   r”   Úadd_fn_target_featureÚadd_fn_asan_attrrd   Úremove_fn_attrÚswap_mirÚdump_mirÚset_all_fn_arg_inregr   rp   rq   rN   Úlink_extern_libsrf   Úneed_extern_libÚlenÚoptimize_moduleÚOPTIMIZE_O3r‘   Úhas_architected_sgprsÚscalarize_packed_fopsÚ#add_scalarize_packed_fops_llvm_passÚcleanup_bitcode_metadataÚdisable_print_inline)rN  rà   r¿   r  r  Ú_HIPBackend__HIP_FTZr  Úllvm_modÚtarget_featuresr[  ÚfnsÚ	kernel_fnÚcluster_dimÚtotal_warps_numÚtotal_num_warpsÚdenormal_moderW   rY   rv   ÚpathsÚpaths   &&&                  r   rS  ÚHIPBackend.make_llir^  sx  € àˆä�_Š_˜SŸ[™[Ó)ˆØ
�‰ÔÜ�
‰
×Ñ×6Ñ6°r¿<¹<ÔHÜ�
‰
×Ñ×7Ñ7¸¿L¹LÔIÜ�‰×$Ñ$ RÔ(Ü�‰× Ñ  Ô$Ü�‰×*Ñ*¨2Ô.ð �w×3Ñ3Ô3Ô8KÈGÏLÉL×8YÒ8YÜ�N‰N×6Ñ6°r¸5ÔAÜ�
‰
×Ñ×5Ñ5°b¿,¹,ÔGà�w×3Ñ3Ô3Ô8KÈGÏLÉL×8YÒ8YÜ�N‰N×4Ñ4°RÔ8Ü�L‰L×*Ñ*¨2Ô.Ü�M‰M×!Ñ! "Ô%Ü�‰×9Ñ9¸"Ô=ä×%×%Ð%Ü×&Ñ&×,Ñ,Ð-@À"ÇkÁkÔRÜ�‰×9Ñ9¸"Ô=ð ˆ	Ü�
‰
×Ñ×(Ñ(¨¯\©\¸9ÔEÜ�
‰
×Ñ×6Ñ6°r¿<¹<ÔHÜ�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!ä�‰×'Ñ'¨Ô+Ü�‰×*Ñ*¨2Ô.Ü�‰×'Ñ'¨Ô+Ü�‰×Ñ˜bÔ!Ü�‰×$Ñ$ RÔ(ô ×%×%Ð%Ü×&Ñ&×,Ñ,Ð-=¸rÇ;Á;ÔOä× Ñ ×2×2Ð2¼5×;LÑ;L×;o×;oÐ;oÜ�M‰M×&Ñ& rÔ*ä�
‰
×Ñ×5Ñ5°b¿,¹,È	ÔRä�‰×5Ñ5°bÔ9Ø
�‰ˆs�KÔ ä×Ñ×?×?Ð?ä×$Ñ$×6×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ˆrÑ4ˆÐPÜÑG¡sÓGÈÓNˆ	ßÜÐ?Ó@Ð@Ø×"Ñ"¤3×#AÑ#AÔBà˜zÕ*ˆØ×ÑÐ3¸°}ÀDÐ5IÔJà!×+Ñ+ˆØ×*Ñ*Ð+@ÓAˆØÒ&Ø-ˆOØ×ÑÐ;¸rÀ/×RcÑRcÕBcÐAdÐ=eÔfØ×ÑÐ7¸Ô@ð ×Ñ 1Ô$Ø×!Ñ!Ð"7¸G×<PÑ<PÐ;QÐQRÐSZ×SgÑSgÐRhÐ9iÔjä& w§|¡|×4Ò4¸×9JÑ9JÈaÔ9OØ×!Ñ!Ð"9¸8ÔDà+2×+E×+EÐ+E™È6ˆØ×ÑÐ4°mÔDä×Ñ×(×(Ð(Ø×+Ñ+¨HÔ5Ø×&Ñ&Ô(Ø"×0Ô0‰KˆD�%Ø×$Ñ$ TÔ*Ø×!Ñ! $¨Ö.ñ 1ð �<‰<˜9Ô$¬e¯i©i×.@×.@Ð.@ÄEÇIÁI×DV×DVÐDVÜ×$Ò$ YÔ/ä×Ñ×(×(Ð(Ü!¤(›^×2Ñ2°UÕ:ˆNä�N \Õ1Ó2Ü�N YÕ.Ó/Ü�N YÕ.Ó/ðˆEô
 ×!Ò! (¨EÕ2Ø× × Ð Ø.5×.AÒ.AÔiÑ.A™l˜t TÄS×EXÒEXÐYaÐcg×Eh—T�TÑ.AˆEÑiÜ�5‹z˜AŒ~Ü×%Ò% h°Ô6ä×Ò˜X¤t×'7Ñ'7¸¿¹ÀrÈ2Èw×OgÑOgØ48õ	:ô ×$Ò$ W§\¡\×2Ò2Ø×$Ñ$Ð%?Ô@Ø×$Ñ$Ð%?Ô@Ø×$Ñ$Ð%?Ô@ä�9‰9×*×*Ð*Ü×3Ò3°IÔ>ð !0�ÑØ ×-Ñ-¨lÓ;ˆ�ÑØ*-×*:Ñ*:Ð;[Ó*\ˆÐ&Ñ'Ø+.×+;Ñ+;Ð<aÓ+bˆÐ'Ñ(Ø+.×+;Ñ+;Ð<]Ó+^×+cÐ+cÐbcˆÐ'Ñ(Ø,/×,<Ñ,<Ð=cÓ,d×,iÐ,iÐhiˆÐ(Ñ)ä×$Ò$ XÔ.ô 	× Ò  Ô*Ü�8‹}Ðùòk Qùól js   Ü%pÝpè7"pépc           
     ód  € \         P                  ! R V 4      p\        V4      ^8X  g   Q hV^ ,          VR&   . p\        VP                  4      '       d   VP                  R4       \        VP                  4      p\        P                  ! V P                  R4      4      P                  4       pV^ ,          R,           V,           p\        P                  ! V \        P                  VP                  WTVP                  V4      p\        P                   ! V \        P                  VP                  WTVP                  V4       \"        P                  P$                  '       d,   \"        P                  P&                  '       g   \)        R4      h\"        P                  P&                  '       d•   \        P*                  ! \,        P.                  P1                  \"        P                  P&                  VR,           4      \        P                  VP                  WTVP                  R\"        P                  P$                  4      p	M>\        P2                  ! V \        P                  VP                  WTVP                  RR4      p	\"        P                  P4                  '       d   \7        R4       \7        V	4       V	# )	z3define amdgpu_kernel void @([a-zA-Z_][a-zA-Z0-9_]*)rW   zamdgpu-expert-scheduling-moder|   rz   zATRITON_SWAP_MIR_ENABLE_MISCHED requires TRITON_SWAP_MIR to be setz.txtFz!// -----// AMDGCN Dump //----- //)ÚreÚfindallr“  r@   r(   rS   rI   r   r€   r�   r‚   r   Útranslate_to_mirr	   r  r‘   Údump_sched_dagr
   Úswap_mir_enable_mischedrŽ  rÑ   Útranslate_mir_to_asmÚosr¥  r}   Útranslate_to_asmÚdump_amdgcnÚprint)
rN  rà   r¿   ÚnamesÚflagsÚfeaturesÚir_hashÚdump_file_idrz   Úamdgcns
   &&&       r   Úmake_amdgcnÚHIPBackend.make_amdgcn  sá  € ô
 —
’
ÐQÐSVÓWˆÜ�5‹z˜QŒÐˆØ  �8ˆ�ÑàˆÜ'¨¯©×5Ò5Ø�L‰LÐ8Ô9Ü.¨w¯|©|Ó<ˆÜ—.’. §¡¨GÓ!4Ó5×?Ñ?ÓAˆØ˜Q•x #•~¨Õ/ˆÜ×!Ò! #¤s×'8Ñ'8¸'¿,¹,ÈÐY`×YqÑYqØ".ó0ˆä×Ò˜C¤×!2Ñ!2°G·L±LÀ(ÐSZ×SkÑSkØ(ô	*ä�9‰9×,×,Ð,´U·Y±Y×5G×5GÐ5GÜÐ`ÓaÐaÜ�9‰9××ÐÜ×.Ò.¬r¯w©w¯|©|¼E¿I¹I×<NÑ<NÐP\Ð_eÕPeÓ/fÜ/2×/@Ñ/@À'Ç,Á,ÐPXØ/6×/GÑ/GÈÔPU×PYÑPY×PqÑPqós‰Fô ×*Ò*¨3´×0AÑ0AÀ7Ç<Á<ÐQYØ+2×+CÑ+CÀUÈEóSˆFä�9‰9× × Ð ÜÐ5Ô6Ü�&ŒMØˆr   c                óŽ  € . p\         P                  P                  '       d   VP                  R 4       \	        VP
                  4      ;p'       d   VP                  V4       \        P                  ! WP
                  RP                  V4      4      p\        P                  ! 4       ;_uu_ 4       p\        P                  ! 4       ;_uu_ 4       p\        VP                  R4      ;_uu_ 4       pVP                  V4       RRR4       \        P                  ! VP                  VP                  4       RRR4       \        VP                  R4      ;_uu_ 4       p	V	P                  4       p
RRR4       RRR4       X
#   + '       g   i     L„; i  + '       g   i     Ld; i  + '       g   i     L@; i  + '       g   i     X
# ; i)rT  rK   ÚwbNÚrb)r
   rr  r}  rS   rI   r(   r	   Úassemble_amdgcnr}   ÚtempfileÚNamedTemporaryFileÚopenrW   ÚwriteÚ
link_hsacoÚread)rN  rà   r¿   r�  Útrue16r´   Útmp_outÚtmp_inÚfd_inÚfd_outr  s   &&&        r   Ú
make_hsacoÚHIPBackend.make_hsaco7  s%  € àˆÜ×Ñ×(×(Ð(Ø×"Ñ" 8Ô,Ü0°·±Ó>Ð>ˆ6Ö>Ø×"Ñ" 6Ô*Ü×#Ò# C¯©°s·x±xÀÓ7PÓQˆÜ×(Ò(×*Ô*¨gÜ×,Ò,×.Ô.°&Ü˜&Ÿ+™+ t×,Ô,°Ø—K‘K Ô&÷ -ä—’˜vŸ{™{¨G¯L©LÔ9÷ /ô �g—l‘l D×)Ô)¨VØ—k‘k“m�÷ *÷ +ð ˆ
÷ -×,ú÷ /×.ú÷ *×)ú÷ +Ö*ð ˆ
úsT   Â/F3ÃF	Ã,E:Ã>3F	Ä1&F3ÅF 	Å(F3Å:F
ÆF	ÆFÆF3Æ F0Æ+F3Æ3G	c                óZ  a a€ V\         P                  8X  d   VV 3R  lVR&   VV 3R lVR&   MV\         P                  8X  d   VV 3R lVR&   VV 3R lVR&   VV 3R lVR&   VV 3R	 lVR
&   \        P                  P
                  e&   \        P                  P                  S VSVR4       R# R# )c                 ó(   <€ SP                  WS4      # rÞ   r  ©rN  rà   r¿   rt   s   &&€€r   r    Ú'HIPBackend.add_stages.<locals>.<lambda>J  s   ø€ °4·>±>À#ÐQXÔ3Yr   r  c                 ó(   <€ SP                  WS4      # rÞ   )r#  rÍ  s   &&€€r   r    rÎ  K  s   ø€ °D·O±OÀCÐSZÔ4[r   Úttgirc                 ó(   <€ SP                  WS4      # rÞ   )rF  rÍ  s   &&€€r   r    rÎ  M  s   ø€ °D×4GÑ4GÈÐW^Ô4_r   c                 ó(   <€ SP                  WS4      # rÞ   )rS  rÍ  s   &&€€r   r    rÎ  N  s   ø€ ¨t¯~©~¸cÈWÔ/Ur   Úllirc                 ó(   <€ SP                  WS4      # rÞ   )r¸  rÍ  s   &&€€r   r    rÎ  O  s   ø€ °×1AÑ1AÀ#ÐQXÔ1Yr   r·  c                 ó(   <€ SP                  WS4      # rÞ   )rÉ  rÍ  s   &&€€r   r    rÎ  P  s   ø€ °·±ÀÈwÔ0Wr   r´   N)r   ÚTRITONÚGLUONr
   rÎ   Úadd_stages_inspection_hook)rt   Ústagesr¿   rÕ   s   f&f&r   Ú
add_stagesÚHIPBackend.add_stagesH  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                ó   € V P                    # rÞ   r"   )rt   s   &r   r…   ÚHIPBackend.hashT  s   € à—+‘+�Ðr   c                ó4   <€ V ^8„  d   Qh/ RS[ ,          ;R&   # )r   Nrõ   )r™   )r   rš   s   "€r   r   r¬   Š   s   ø‡ ‚ ðL ™T•kÑ(òM r   )r·   )!r¢   r£   r¤   r¥   rï   Ú%supports_native_tensor_specializationÚstaticmethodr¯   r¶   rÀ   rÛ   rá   rå   rë   rî   rõ   rý   r  r  r  r#  rF  rS  r¸  rÉ  rÚ  Ú	functoolsÚ	lru_cacher…   r¦   r§   r¨   Ú__classcell__)r¸   rš   s   @@r   rª   rª   Š   sE  ù‡ € Ø€OØ,1Ð)à÷'ó ð'÷"ó "÷
%ð %÷"ð "ò<
ò?÷>ð >ò
:ð %)Ðàñó ðð" ñó ðð ñó ðð ñó ðð  ñ@ó ð@ðD ñó ðð. ñuó ðuðn ñó ðð@ ñó ðò 
\ð ×ÒÓñ ó ð ÷W … r   rª   )(Útriton.backends.compilerr   r   r   Útriton._C.libtritonr   r   r   r	   Útritonr
   Údataclassesr   Útypingr   r   r   Útypesr   r®  r   r¾  r¨  rá  rn   Úpathlibr   r#   r*   r1   r4   r8   r=   r@   rB   rD   rI   rZ   r]   rª   r   r   r   Ú<module>rë     sœ   ðß EÑ Eß 5Ó 5Ý Ý !ß #Ñ #Ý Û 	Û Û Û 	Û Û Ý õ0òRò
\ò
mòòòò3òò>ò	ñ �$Ô÷:?ð :?ó ð:?ôzL �ö L r   