Ë
    ÏÍ:jT  ã                  óæ  — d dl mZ d dlZd dlZddlmZmZ ddlmZ ddlm	Z	 ddlm
Z
mZ ddlmZmZ dd	lmZ dd
lmZmZmZ ddlmZ ddl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Z de iZ!dZ"de"iZ#d„ Z$ G d„ d«      Z% G d„ d«      Z& ejN                  «       d„ «       Z( ejN                  «       d„ «       Z)d„ Z*d$d„Z+ G d„ d«      Z,d%d„Z-d&d„Z. G d„ d«      Z/ G d „ d!e0«      Z1 G d"„ d#«      Z2y)'é    )ÚannotationsNé   )Úget_cache_invalidating_env_varsÚir)Úbackends)ÚLanguage)ÚBaseBackendÚ	GPUTarget)Ú__version__Úknobs)ÚOutOfResources)Úget_cache_managerÚget_dump_managerÚget_override_manager)Údriver)Úget_sass)ÚPathz=\.(?:visible|extern)\s+\.(?:entry|func)\s+(\w+)\s*\(([^)]*)\)Úptxz\.param\s+\.(\w+)c                óÐ   — t        j                  d| «      }t        j                  d| «      }|�yt        j                  dd| «      } |�dt        |j	                  d«      «      z   S | S )Nz!tt\.ptr<([^,]+)ztt.nv_tma_desc = 1Ú	nvTmaDescz {[^}]+}Ú Ú*é   )ÚreÚsearchÚsubÚconvert_type_reprÚgroup)ÚxÚmatchÚtmas      úm/home/mcse/projects/srt_converter/srt-converter-venv/lib/python3.12/site-packages/triton/compiler/compiler.pyr   r   '   sd   € ô �I‰IÐ)¨1Ó-€EÜ
�)‰)Ð)¨1Ó
-€CØ
€ØÜ
�‰ˆ{˜B Ó"€AØÐØÔ& u§{¡{°1£~Ó6Ñ6Ð6Ø€Hó    c                  ó(   — e Zd Zddd„Zd„ Zd„ Zd„ Zy)Ú	ASTSourceNc                óÜ  — || _         t        j                  | _        d| _        |j
                  | _        || _        t        «       | _	        |�g|j                  «       D ]T  \  }}t        |t        «      r|j                  j                  |«      fn|}t        |t        «      sJ ‚|| j                  |<   ŒV |xs
 t        «       | _        t        | j                  t        «      rLt#        | j                  j%                  d«      «      D ��ci c]  \  }}||j'                  «       “Œ c}}| _        y | j                  j)                  «       D ]  }t        |t        «      rŒt+        d«      ‚ y c c}}w )NÚttirú,zSignature keys must be string)Úfnr   ÚTRITONÚlanguageÚextÚ__name__ÚnameÚ	signatureÚdictÚ	constantsÚitemsÚ
isinstanceÚstrÚ	arg_namesÚindexÚtupleÚattrsÚ	enumerateÚsplitÚstripÚkeysÚ	TypeError)Úselfr)   r/   Ú
constexprsr8   ÚkÚvs          r"   Ú__init__zASTSource.__init__6   s$  € ØˆŒÜ Ÿ™ˆŒØˆŒØ—K‘KˆŒ	Ø"ˆŒÜ›ˆŒØÐ!Ø"×(Ñ(Ó*ò &‘��1Ü1;¸A¼sÔ1C�R—\‘\×'Ñ'¨Ó*Ñ-È�Ü! !¤UÔ+Ð+Ð+Ø$%�—‘˜qÒ!ð&ð ’_œd›fˆŒ
Ü�d—n‘n¤cÔ*Ü7@ÀÇÁ×AUÑAUÐVYÓAZÓ7[×\©t¨q°!˜a §¡£™lÓ\ˆD�Nà—^‘^×(Ñ(Ó*ò E�Ü! !¤SÕ)Ü#Ð$CÓDÐDñEùó ]s   ÄE(c           	     óì  — t        | j                  j                  «       «      D ��cg c]  \  }}|‘Œ	 }}}d„ }dj                  t        | j                  j                  «       «      D ��cg c]  \  }} ||«      ‘Œ c}}«      }| j
                  j                  › dt        | j                  «      › d|› d|› �}t        j                  |j                  d«      «      j                  «       S c c}}w c c}}w )Nc                óH   — t        | d«      r| j                  S t        | «      S )NÚ	cache_key)ÚhasattrrE   r4   )r   s    r"   ú<lambda>z ASTSource.hash.<locals>.<lambda>L   s   € ¬7°1°kÔ+B˜AŸK™K€ ÌÈAË€ r#   ú-úutf-8)Úsortedr/   r2   Újoinr1   r)   rE   r4   r8   ÚhashlibÚsha256ÚencodeÚ	hexdigest)r>   r@   rA   Ú
sorted_sigÚget_keyÚconstants_keyÚkeys          r"   ÚhashzASTSource.hashJ   s½   € Ü$*¨4¯>©>×+?Ñ+?Ó+AÓ$B×C™D˜A˜q’aÐCˆ
ÑCÙNˆØŸ™¼ÀÇÁ×@TÑ@TÓ@VÓ9W×!X±°°A¡'¨!¥*Ó!XÓYˆØ—‘×"Ñ"Ð# 1¤S¨¯©£_Ð$5°Q°z°lÀ!ÀMÀ?ÐSˆÜ�~‰~˜cŸj™j¨Ó1Ó2×<Ñ<Ó>Ð>ùó	 Dùã!Xs   §C*Á*C0
c                ó>   — ddl m}  || j                  | ||||¬«      S )Nr   )Úast_to_ttir)ÚcontextÚoptionsÚcodegen_fnsÚ
module_map)Úcode_generatorrV   r)   )r>   rX   rY   rZ   rW   rV   s         r"   Úmake_irzASTSource.make_irQ   s$   € Ý/Ù˜4Ÿ7™7 D°'À7ÐXcØ&0ô2ð 	2r#   c                ó   — t        «       S ©N)r0   ©r>   s    r"   Úparse_optionszASTSource.parse_optionsV   s	   € Ü‹vˆr#   ©NN©ÚreturnÚNone©r-   Ú
__module__Ú__qualname__rB   rT   r\   r`   © r#   r"   r%   r%   4   s   „ ôEò(?ò2ó
r#   r%   c                  ó$   — e Zd Zd„ Zd„ Zd„ Zd„ Zy)ÚIRSourcec                óä  — || _         t        |«      }|j                  dd  | _        t        j
                  | _        |j                  «       | _        t        j                  |«       |j                  |«       | j                  dk(  r¼t        j                  t        | j                     | j                  t        j                  «      }|j                  d«      | _        |j                  d«      }t        j"                  t$        | j                     |«      }t'        |«      D ��ci c]  \  }}|t)        |«      “Œ c}}| _        y t        j,                  | j                   |«      | _        | j.                  j1                  «       }	d|	z   | _        | j.                  j3                  |	«      }
| j.                  j5                  |
«      }t'        |«      D ��ci c]  \  }}||“Œ
 c}}| _        y c c}}w c c}}w )Nr   r   r   ú@)Úpathr   Úsuffixr,   r   r*   r+   Ú	read_textÚsrcr   Úload_dialectsr   r   Úprototype_patternÚ	MULTILINEr   r.   ÚfindallÚarg_type_patternr9   r   r/   Úparse_mlir_moduleÚmoduleÚget_entry_func_nameÚget_functionÚget_function_signature)r>   rm   rW   Úbackendr    r/   Útypesr@   ÚtyÚfn_nameÚfuncOpÚfunc_tys               r"   rB   zIRSource.__init__\   sn  € ØˆŒ	Ü�D‹zˆØ—;‘;˜q˜r�?ˆŒÜ Ÿ™ˆŒØ—>‘>Ó#ˆŒÜ
×Ñ˜Ô!Ø×Ñ˜gÔ&ð �8‰8�uÒÜ—I‘IÔ/°·±Ñ9¸4¿8¹8ÄRÇ\Á\ÓRˆEØŸ™ A›ˆDŒIØŸ™ A›ˆIÜ—J‘JÔ/°·±Ñ9¸9ÓEˆEÜDMÈeÓDT×U¹5¸1¸b˜aÔ!2°2Ó!6Ñ6ÓUˆD�Nä×.Ñ.¨t¯y©y¸'ÓBˆDŒKØ—k‘k×5Ñ5Ó7ˆGØ˜g™ˆDŒIØ—[‘[×-Ñ-¨gÓ6ˆFØ—k‘k×8Ñ8¸Ó@ˆGÜ1:¸7Ó1C×D©¨¨2˜a ™eÓDˆD�Nùó Vùó Es   Ä"G&ÇG,c                óz   — t        j                  | j                  j                  d«      «      j	                  «       S )NrI   )rL   rM   rp   rN   rO   r_   s    r"   rT   zIRSource.hashu   s'   € Ü�~‰~˜dŸh™hŸo™o¨gÓ6Ó7×AÑAÓCÐCr#   c                ó<   — || j                   _        | j                   S r^   )rw   rW   )r>   rX   rY   rZ   rW   s        r"   r\   zIRSource.make_irx   s   € Ø%ˆ�‰ÔØ�{‰{Ðr#   c                ó„   — | j                   dk(  r(| j                  j                  d«      }|€J d«       ‚d|iS t        «       S )NÚttgirzttg.num-warpsz'Unable to parse ttg.num-warps attributeÚ	num_warps)r,   rw   Úget_int_attrr0   )r>   r…   s     r"   r`   zIRSource.parse_options|   sF   € Ø�8‰8�wÒØŸ™×0Ñ0°ÓAˆIØÐ(ÐSÐ*SÓSÐ(Ø Ð+Ð+Ü‹vˆr#   Nre   rh   r#   r"   rj   rj   Z   s   „ òEò2Dòór#   rj   c                 óŠ  — dd l } t        j                  j                  t        j                  j                  t        j                  j	                  t
        «      «      «      }g }t        t
        d«      5 }|t        j                  |j                  «       «      j                  «       gz  }d d d «       t        j                  j                  |d«      dft        j                  j                  |d«      dfg}|D ]•  \  }}| j                  |g|¬«      D ]y  }t        |j                  j                  |j                  «      j                   d«      5 }|t        j                  |j                  «       «      j                  «       gz  }d d d «       Œ{ Œ— t        j                  «       }t#        j$                  d«      j'                  d	«      d
   }	t        t        j                  j                  |dd|	› �«      d«      5 }	 |j                  d«      }
|
sn|j)                  |
«       Œ&	 d d d «       |j+                  |j                  «       «       t        j                  j                  |d«      }| j                  |gd¬«      D ]y  }t        |j                  j                  |j                  «      j                   d«      5 }|t        j                  |j                  «       «      j                  «       gz  }d d d «       Œ{ t,        › dj                  |«      z   S # 1 sw Y   �ŒjxY w# 1 sw Y   �ŒxY w# 1 sw Y   �ŒxY w# 1 sw Y   ŒÅxY w)Nr   ÚrbÚcompilerztriton.compiler.r   ztriton.backends.)ÚprefixÚ
EXT_SUFFIXú.éÿÿÿÿÚ_Cz
libtriton.i   r+   ztriton.language.rH   )ÚpkgutilÚosrm   ÚdirnameÚabspathÚ__file__ÚopenrL   rM   ÚreadrO   rK   Úwalk_packagesÚmodule_finderÚ	find_specr.   ÚoriginÚ	sysconfigÚget_config_varr:   ÚupdateÚappendr   )r�   ÚTRITON_PATHÚcontentsÚfÚpath_prefixesrm   rŠ   ÚlibÚlibtriton_hashr,   ÚchunkÚlanguage_paths               r"   Ú
triton_keyr¦   „   sº  € ãÜ—'‘'—/‘/¤"§'¡'§/¡/´"·'±'·/±/Ä(Ó2KÓ"LÓM€KØ€Hä	Œh˜Ó	ð ; Ø”W—^‘^ A§F¡F£HÓ-×7Ñ7Ó9Ð:Ñ:ˆ÷;ô 
�‰�‰�k :Ó	.Ð0BÐCÜ	�‰�‰�k :Ó	.Ð0BÐCð€Mð &ò C‰ˆˆfØ×(Ñ(¨$¨¸Ð(Ó?ò 	CˆCÜ�c×'Ñ'×1Ñ1°#·(±(Ó;×BÑBÀDÓIð CÈQØœWŸ^™^¨A¯F©F«HÓ5×?Ñ?ÓAÐBÑB�÷Cð Cñ	CðCô —^‘^Ó%€NÜ
×
"Ñ
" <Ó
0×
6Ñ
6°sÓ
;¸BÑ
?€CÜ	Œb�g‰g�l‰l˜;¨°
¸3¸%Ð.@ÓAÀ4Ó	Hð )ÈAØØ—F‘F˜7“OˆEÙØØ×!Ñ! %Ô(ð	 ð ÷	)ð ‡O�O�N×,Ñ,Ó.Ô/ä—G‘G—L‘L ¨jÓ9€MØ×$Ñ$ m _Ð=OÐ$ÓPò ?ˆÜ�#×#Ñ#×-Ñ-¨c¯h©hÓ7×>Ñ>ÀÓEð 	?ÈØœŸ™¨¯©«Ó1×;Ñ;Ó=Ð>Ñ>ˆH÷	?ð 	?ð?ô ˆ]˜cŸh™h xÓ0Ñ0Ð0÷7;ñ ;ú÷Cñ Cú÷)ñ )ú÷	?ð 	?ús0   Á46LÅ6LÇ9(L,Ê96L9ÌLÌL)Ì,L6Ì9M	c                óZ   — t         j                  j                  j                  | «      d   S )NÚmax_shared_mem)r   ÚactiveÚutilsÚget_device_properties)Údevices    r"   r¨   r¨   ¨   s#   € ä�=‰=×Ñ×4Ñ4°VÓ<Ð=MÑNÐNr#   c                óì   — |dk(  s|dk(  rt        j                  | |«      }||_        |S |dk(  s
|dk(  s|dk(  rt        | «      j	                  «       S |dk(  s|dk(  rt        | «      j                  «       S y )Nr'   r„   Úllirr   ÚamdgcnÚcubinÚhsaco)r   rv   rW   r   ro   Ú
read_bytes)Ú	full_namer,   rW   rw   s       r"   Úparser´   ­   sx   € Ø
ˆf‚}˜˜wšÜ×%Ñ% i°Ó9ˆØ ˆŒØˆØ
ˆf‚}˜˜uš¨¨xªÜ�I‹×(Ñ(Ó*Ð*Ø
ˆg‚~˜ šÜ�I‹×)Ñ)Ó+Ð+ð (r#   c                ó,  ‡— t         j                  j                  ry| j                  �t	        | j                  «       | j
                  �t	        | j
                  «       ddg}|D �cg c]"  }|j                  dt        j                  «      ‘Œ$ }}| j                  Šg }‰�4t        ˆfd„|D «       «      s|j                  ‰«       ‰j                  Š‰�Œ4t        ||dd «      D ]  \  }}||_        Œ |sd| _	        yd|d   _        |d   | _	        yc c}w )	zŒ
    Removes code_generator.py and related files from tracebacks.

    These are uninteresting to the user -- "just show me *my* code!"
    Nz"/triton/compiler/code_generator.pyz/ast.pyú/c              3  ó‚   •K  — | ]6  }‰j                   j                  j                  j                  |«      sŒ3|–— Œ8 y ­wr^   )Útb_frameÚf_codeÚco_filenameÚendswith)Ú.0r    Útbs     €r"   ú	<genexpr>z#filter_traceback.<locals>.<genexpr>Ð   s/   øè ø€ ÒV˜¨2¯;©;×+=Ñ+=×+IÑ+I×+RÑ+RÐSTÕ+U”1ÑVùs   ƒ4?¸?r   r�   r   )r   ÚcompilationÚfront_end_debuggingÚ	__cause__Úfilter_tracebackÚ__context__Úreplacer�   ÚsepÚ__traceback__Úanyr�   Útb_nextÚzip)ÚeÚ	BAD_FILESÚbad_fileÚframesÚ	cur_frameÚ
next_framer½   s         @r"   rÂ   rÂ   ¸   s  ø€ ô ×Ñ×,Ò,Øà‡{�{ÐÜ˜Ÿ™Ô%Ø‡}�}Ð Ü˜Ÿ™Ô'ð 	-Øð€Ið @IÖI°8�×!Ñ! #¤r§v¡vÕ.ÐI€IÐIà	
�‰€BØ€FØ
ˆ.ÜÓV˜iÔVÔVØ�M‰M˜"ÔØ�Z‰Zˆð ‰.ô
 $' v¨v°a°b¨zÓ#:ò 'Ñˆ�JØ&ˆ	Õð'ñ Øˆ�à!ˆˆr‰
ÔØ  ™)ˆ�ùò! Js   Á''Dc                  ó,   — e Zd Zdd„Zdd„Zdd„Zdd„Zy)	ÚCompileTimerc                ó`   — t        j                   «       | _        d | _        g | _        d | _        y r^   )ÚtimeÚstartÚir_initialization_endÚlowering_stage_endsÚstore_results_endr_   s    r"   rB   zCompileTimer.__init__à   s'   € Ü ŸI™I›KˆŒ
Ø37ˆÔ"Ø<>ˆÔ Ø/3ˆÕr#   c                ó6   — t        j                   «       | _        y r^   )rÓ   rÕ   r_   s    r"   Úfinished_ir_initializationz'CompileTimer.finished_ir_initializationæ   s   € Ü%)§Y¡Y£[ˆÕ"r#   c                ób   — | j                   j                  |t        j                  «       f«       y r^   )rÖ   r�   rÓ   )r>   Ú
stage_names     r"   Ústage_finishedzCompileTimer.stage_finishedé   s!   € Ø× Ñ ×'Ñ'¨´T·Y±Y³[Ð(AÕBr#   c                ól  — t        j                   «       }| j                  €|| _        n|| _        dd„}g }| j                  }| j                  D ]!  \  }}|j	                  | |||«      f«       |}Œ# t        j                   || j                  | j                  «      | ||| j                  «      ¬«      S )Nc                ó*   — |€yt        || z
  dz  «      S )Nr   i@B )Úint)rÔ   Úends     r"   ÚdeltazCompileTimer.end.<locals>.deltaó   s   € Øˆ{ØÜ˜˜e™ wÑ.Ó/Ð/r#   )Úir_initializationÚlowering_stagesÚstore_results)rÔ   Úfloatrà   zfloat | Nonerc   rß   )rÓ   rÕ   r×   rÖ   r�   r   ÚCompileTimesrÔ   )r>   Ú	timestamprá   Úlowering_stage_durationsÚstage_startrÛ   Ú	stage_ends          r"   rà   zCompileTimer.endì   s¶   € Ü—I‘I“Kˆ	Ø×%Ñ%Ð-Ø)2ˆDÕ&à%.ˆDÔ"ó	0ð
 $&Ð Ø×0Ñ0ˆØ%)×%=Ñ%=ò 	$Ñ!ˆJ˜	Ø$×+Ñ+¨Z¹¸{ÈIÓ9VÐ,WÔXØ#‰Kð	$ô ×!Ñ!Ù# D§J¡J°×0JÑ0JÓKØ4Ù ¨T×-CÑ-CÓDô
ð 	
r#   Nrb   )rÛ   r4   rc   rd   )rc   zknobs.CompileTimes)r-   rf   rg   rB   rÙ   rÜ   rà   rh   r#   r"   rÑ   rÑ   Þ   s   „ ó4ó1óCô
r#   rÑ   c                óB  — t         j                  j                  }|r
t        «       }|€t        j
                  j                  «       }t        |t        «      sJ d«       ‚t        |«      }t        | t        «       }|r8t        | t        «      sJ d«       ‚t        j                  «       }t        | ||«      } | j                  «       }|j                  t!        |xs
 t!        «       fi |¤Ž«      }t#        «       }	t%        «       › d| j'                  «       › d|j'                  «       › d|j'                  «       › dt        t)        |	j+                  «       «      «      › �	}
t-        j.                  |
j1                  d«      «      j3                  «       }t5        |«      }t         j                  j6                  }t         j                  j8                  }t         j                  j:                  }|rt=        | j'                  «       «      nd }|rt?        | j'                  «       «      nd }| j@                  d d }|› d�}|jC                  |«      xs i }|jE                  |«      }t         j                  jF                  }|sF|�DtI        | ||«      }|r3 || |jJ                  jM                  «       |jO                  «       d¬«       |S ||d	œ|jP                  ¥|	¥}tR        |d
<   t!        «       }|jU                  ||| jV                  «       tY        |j[                  «       «      j]                  | j^                  «      }|r|dz  }t        | t        «      s:t        j                  «       }t        j`                  |«       |ja                  |«       |jc                  |«      }|je                  «       }	 | jg                  |||«      }|r'|› d| j^                  › �}|jm                  ||«      ||<   n|› d�}|jm                  ||«      ||<   t         j                  jn                  } |r5| r3|jq                  | jr                  «       tu        d| jr                  › �«       |rjw                  «        tY        |j+                  «       «      |d  D ]õ  \  }!}" |"||«      }#|› d|!› �}|€6|jE                  dd «      x}$rP|$jy                  d|!› �«      r<t{        |$|!|«      }#n.|j}                  |«      x}%rtu        d|%› �«       t{        |%|!|«      }#|r|!dv r|jm                  |#|«      ||<   |�|jm                  |#|«       | |!k(  r0|j}                  |«      }&|#jq                  |&«       tu        d|&› �«       |#}|sŒåj                  |!«       Œ÷ |jm                  t�        j‚                  |t„        ¬«      |d¬«      ||<   |j‡                  ||«       t         j                  jˆ                  s|j‹                  «        |r || ||jO                  «       d¬«       tI        | ||«      S # th        $ r}tk        |«       ‚ d }~ww xY w)Nz target must be of GPUTarget typez'source must be either AST or a filepathrH   rI   é–   ú.jsonT)rp   ÚmetadataÚmetadata_groupÚtimesÚ	cache_hit)rT   ÚtargetÚtriton_versionr   rŒ   z.sourcezCreating new locations for Úir_overridez
Overriding kernel with file )r°   r±   Újson)ÚdefaultF)Úbinary)Fr   r¿   ÚlistenerrÑ   r   r©   Úget_current_targetr3   r
   Úmake_backendr%   r4   r   rW   rj   r`   r0   r   r¦   rT   rJ   r2   rL   rM   rN   rO   r   ÚoverrideÚdump_irÚstore_binary_onlyr   r   r.   Ú	get_groupÚgetÚalways_compileÚCompiledKernelrî   Ú_asdictrà   Ú__dict__r   Ú
add_stagesr+   Úlistr<   r6   r,   rq   Úget_codegen_implementationÚget_module_mapr\   Ú	ExceptionrÂ   ÚputÚ
use_ir_locÚcreate_location_snapshotrm   ÚprintrÙ   r»   r´   Úget_filerÜ   rõ   ÚdumpsÚvarsÚ	put_groupÚenable_asanÚdisable_multithreading)'rp   rò   rX   Úcompilation_listenerÚtimerr{   Ú	ir_sourcerW   Úextra_optionsÚenv_varsrS   rT   Úfn_cache_managerÚenable_overrideÚenable_ir_dumpÚstore_only_binaryÚfn_override_managerÚfn_dump_managerÚ	file_nameÚmetadata_filenamerï   Úmetadata_pathr   Úresrî   ÚstagesÚfirst_stagerY   rZ   rw   rÊ   Úir_filenamer
  r,   Ú
compile_irÚnext_modulerô   r³   Úir_full_names'                                          r"   Úcompiler(    s»  € Ü ×,Ñ,×5Ñ5ÐÙÜ“ˆà€~Ü—‘×1Ñ1Ó3ˆÜ�fœiÔ(ÐLÐ*LÓLÐ(Ü˜6Ó"€GÜ˜s¤IÓ.Ð.€IáÜ˜#œsÔ#ÐNÐ%NÓNÐ#Ü—*‘*“,ˆÜ�s˜G WÓ-ˆà×%Ñ%Ó'€MØ×#Ñ#¤D¨Ò):´D³FÑ$L¸mÑ$LÓM€Gä.Ó0€HÜ‹\ˆN˜!˜CŸH™H›J˜< q¨¯©«Ð(8¸¸'¿,¹,».Ð9IÈÌ3ÌvÐV^×VdÑVdÓVfÓOgÓKhÐJiÐ
j€CÜ�>‰>˜#Ÿ*™* WÓ-Ó.×8Ñ8Ó:€DÜ(¨Ó.Ðô ×'Ñ'×0Ñ0€OÜ×&Ñ&×.Ñ.€NÜ×)Ñ)×;Ñ;ÐÙ>MÔ.¨s¯x©x«zÔ:ÐSWÐÙ6DÔ& s§x¡x£zÔ2È$€Oð
 —‘˜˜#�€IØ$˜+ UÐ+ÐØ%×/Ñ/Ð0AÓBÒHÀb€NØ"×&Ñ&Ð'8Ó9€MÜ×&Ñ&×5Ñ5€NÙ˜mÐ7ä˜S .°$Ó7ˆÙÙ ØØŸ™×-Ñ-Ó/Ø-Ø—i‘i“kØõð ˆ
ð Øñð ×
Ñ
ðð ð	€Hô "-€HÐÑä‹V€FØ×Ñ�v˜w¨¯©Ô5Ü�v—{‘{“}Ó%×+Ñ+¨C¯G©GÓ4€KáØ�qÑˆô �cœ8Ô$Ü—*‘*“,ˆÜ
×Ñ˜Ô!Ø×Ñ˜gÔ&à×4Ñ4°WÓ=€KØ×'Ñ'Ó)€JðØ—‘˜W k°:¸wÓGˆñ
 Ø"˜ 1 S§W¡W IÐ.ˆØ&6×&:Ñ&:¸6À;Ó&Oˆ�{Ò#à"˜ 7Ð+ˆØ&6×&:Ñ&:¸6À;Ó&Oˆ�{Ñ#ä×"Ñ"×-Ñ-€JÙ‘ZØ×'Ñ'¨¯©Ô1ÜÐ+¨C¯H©H¨:Ð6Ô7áØ×(Ñ(Ô*Ü §¡£Ó/°°Ð=ò &‰ˆˆZÙ  ¨Ó2ˆØ"˜ 1 S EÐ*ˆØÐ&ð  (Ÿ|™|¨M¸4Ó@Ð@�Ð@Àk×FZÑFZÐ]^Ð_bÐ^cÐ[dÔFeÜ# K°°gÓ>‘Ø-×6Ñ6°{ÓCÐCˆYÐCÜÐ2°9°+Ð>Ô?Ü 	¨3°Ó8ˆKá! sÐ.HÑ'HØ*:×*>Ñ*>¸{ÈKÓ*XˆN˜;Ñ'ØÐ&Ø×Ñ ¨[Ô9à˜ÒØ+×4Ñ4°[ÓAˆLØ×0Ñ0°Ô>ÜÐ/°¨~Ð>Ô?ØˆÚØ× Ñ  Õ%ð/&ð2 )9×(<Ñ(<¼T¿Z¹ZÈÔZ^Ô=_ÐarØDIð )=ó )K€NÐ$Ñ%à×ÑÐ0°.ÔAô ×Ñ×(Ò(Ø×&Ñ&Ô(ñ Ù ¨xÈÐ^c×^gÑ^gÓ^iØ',õ	.ô ˜#˜~¨tÓ4Ð4øô ò Ü˜ÔØûðús   Î X Ø	XØXØXc                ó  — t        j                  «       D �cg c]*  }|j                  j                  | «      sŒ|j                  ‘Œ, }}t	        |«      dk7  r't        t	        |«      › d| j                  › d|› d�«      ‚ |d   | «      S c c}w )Nr   z! compatible backends for target (z) (z). There should only be one.r   )r   Úvaluesr‰   Úsupports_targetÚlenÚRuntimeErrorr{   )rò   r   Úactivess      r"   rú   rú   –  s‡   € Ü#+§?¡?Ó#4Ö[˜a¸¿
¹
×8RÑ8RÐSYÕ8Zˆq�z‹zÐ[€GÐ[Ü
ˆ7ƒ|�qÒÜÜ�7‹|ˆnÐ=¸f¿n¹nÐ=MÈSÐQXÐPYÐYuÐvóxð 	xàˆ7�1‰:�fÓÐùò	 \s
   — B¸Bc                  ó   — e Zd Zd„ Zd„ Zd„ Zy)ÚLazyDictc                ó    — || _         g | _        y r^   )ÚdataÚextras)r>   r2  s     r"   rB   zLazyDict.__init__   s   € ØˆŒ	Øˆ�r#   c                ó¤   — | j                   D ]  \  }}| j                   ||Ž z  | _        Œ | j                   j                  «        | j                  S r^   )r3  r2  Úclear©r>   ÚfuncÚargss      r"   rÿ   zLazyDict.get¤  sG   € ØŸ+™+ò 	0‰JˆD�$ØŸ	™	¡D¨$ KÑ/ˆD�Ið	0à�‰×ÑÔØ�y‰yÐr#   c                ó>   — | j                   j                  ||f«       y r^   )r3  r�   r6  s      r"   ÚaddzLazyDict.addª  s   € Ø�‰×Ñ˜D $˜<Õ(r#   N)r-   rf   rg   rB   rÿ   r:  rh   r#   r"   r0  r0  ž  s   „ òòó)r#   r0  c                  ó   — e Zd Zd„ Zy)ÚAsmDictc                óT   — |dk(  rt        | d   «      }nt        d|z  «      ‚|| |<   |S )NÚsassr°   zUnknown key: '%s')r   ÚKeyError)r>   rS   Úvalues      r"   Ú__missing__zAsmDict.__missing__°  s6   € à�&Š=Ü˜T '™]Ó+‰EäÐ.°Ñ4Ó5Ð5àˆˆS‰	Øˆr#   N)r-   rf   rg   rA  rh   r#   r"   r<  r<  ®  s   „ ór#   r<  c                  ó6   ‡ — e Zd Zd„ Zd„ Zˆ fd„Zd„ Zd„ Zˆ xZS )r  c           	     óÔ  — ddl m} t        d„ |j                  «       D «       «      }t	        j
                  |j                  «       «      }t        |d   «      |d<   |d   }t        |d   |d   |d   «      |d<    |d	t        t        |j                  «       «      «      «      } |di |¤Ž| _        t        | j                  j                  «      }	|	j                  | j                  «      | _        || _        || _        | j                  j&                  | _        |j                  «       D �
�cg c]"  \  }
}|
j)                  d
«      rŒt+        |«      ‘Œ$ }}
}|	j,                  }t/        |D �ci c]B  }|j0                  dd  |j0                  dd  |k(  r|j3                  «       n|j                  «       “ŒD c}«      | _        | j4                  |   | _        d | _        d | _        y c c}}
w c c}w )Nr   )Ú
namedtuplec              3  ó\   K  — | ]$  \  }}|j                  d «      sŒt        |«      –— Œ& y­w)rí   N)r»   r   )r¼   ÚcÚps      r"   r¾   z*CompiledKernel.__init__.<locals>.<genexpr>¿  s$   è ø€ Ò`©$¨!¨QÈAÏJÉJÐW^ÕL_œd 1ŸgÑ`ùs   ‚,œ,Úcluster_dimsrò   r{   ÚarchÚ	warp_sizeÚKernelMetadatarí   r   rh   )ÚcollectionsrD  Únextr2   rõ   Úloadsro   r7   r
   rJ   r  r<   rî   rú   rò   Úpack_metadataÚpacked_metadatarp   rT   r.   r»   r   Ú
binary_extr<  rn   r²   ÚasmÚkernelrw   Úfunction)r>   rp   rï   rT   rD  r   rî   rò   rK  r{   rF  rG  Ú	asm_filesrQ  Úfiles                  r"   rB   zCompiledKernel.__init__½  s¨  € Ý*ÜÑ`°.×2FÑ2FÓ2HÔ`ÓaˆÜ—:‘:˜m×5Ñ5Ó7Ó8ˆÜ#(¨°.Ñ)AÓ#Bˆ�Ñ à˜(Ñ#ˆÜ& v¨iÑ'8¸&À¹.È&ÐQ\ÑJ]Ó^ˆ�ÑÙ#Ð$4´f¼TÀ(Ç-Á-Ã/Ó=RÓ6SÓTˆÙ&Ñ2¨Ñ2ˆŒÜ˜tŸ}™}×3Ñ3Ó4ˆØ&×4Ñ4°T·]±]ÓCˆÔØˆŒØˆŒ	Ø—M‘M×&Ñ&ˆŒ	à)7×)=Ñ)=Ó)?×[¡  AÀqÇzÁzÐRYÕGZ”T˜!•WÐ[ˆ	Ñ[Ø×'Ñ'ˆ
Üà!ö
àð �K‰K˜˜ˆO°$·+±+¸a¸b°/ÀZÒ2O˜TŸ_™_Ô.ÐUY×UcÑUcÓUeÑeò
ó ˆŒð —h‘h˜zÑ*ˆŒð ˆŒØˆ�ùó \ùò
s   Ä*GÅGÅ)AG%c                ó  — | j                   �y t        j                  j                  «       }t        j                  j	                  | j
                  | j                  «      | _        t        |«      }| j                  j                  |kD  r!t        | j                  j                  |d«      ‚t        | j                  d«      rR| j                  j                  �<d}| j                  j                  |kD  r!t        | j                  j                  |d«      ‚t        j                  j                  j                  | j                  | j                   | j                  j                  |«      \  | _         | _        | _        | _        | _        t        j                  j+                  «       j,                  }| j                  j.                  |z  | j(                  kD  r.t        | j                  j.                  |z  | j(                  d«      ‚y )Nzshared memoryÚ	tmem_sizei   ztensor memoryÚthreads)rw   r   r©   Úget_current_deviceÚlauncher_clsrp   rî   Úrunr¨   Úsharedr   rF   rX  rª   Úload_binaryr.   rS  rT  Ún_regsÚn_spillsÚn_max_threadsrù   rJ  r…   )r>   r¬   Ú
max_sharedÚmax_tmem_sizerJ  s        r"   Ú_init_handleszCompiledKernel._init_handlesÚ  sr  € Ø�;‰;Ð"ØÜ—‘×1Ñ1Ó3ˆä—=‘=×-Ñ-¨d¯h©h¸¿¹ÓFˆŒä# FÓ+ˆ
Ø�=‰=×Ñ *Ò,Ü  §¡×!5Ñ!5°zÀ?ÓSÐSÜ�4—=‘= +Ô.°4·=±=×3JÑ3JÐ3VàˆMØ�}‰}×&Ñ&¨Ò6Ü$ T§]¡]×%<Ñ%<¸mÈ_Ó]Ð]äU[×UbÑUb×UhÑUh×UtÑUtØ�I‰I�t—{‘{ D§M¡M×$8Ñ$8¸&óVBÑRˆŒ�T”] D¤K°´ÀÔ@Rä—M‘M×4Ñ4Ó6×@Ñ@ˆ	Ø�=‰=×"Ñ" YÑ.°×1CÑ1CÒCÜ  §¡×!8Ñ!8¸9Ñ!DÀd×FXÑFXÐZcÓdÐdð Dr#   c                óL   •— |dk(  r| j                  «        t        ‰| �	  |«      S )Nr\  )rd  ÚsuperÚ__getattribute__)r>   r.   Ú	__class__s     €r"   rg  zCompiledKernel.__getattribute__ð  s&   ø€ Ø�5Š=Ø×ÑÔ Ü‰wÑ'¨Ó-Ð-r#   c                óø  — t         j                  j                  €y t        | j                  | j
                  |dœ«      }t        | j                  t        «      r | j                  j                  j                  €|S i }d}t        | j                  j                  j                  «      D ]  \  }}||   ||<   |dz  }Œ |j                  | j                  j                  j                  || j                  |f«       |S )N)r.   rT  Ústreamr   r   )r   ÚruntimeÚlaunch_enter_hookr0  r.   rT  r3   rp   r%   r)   Úlaunch_metadatar9   r5   r:  rî   )	r>   Úgridrj  r8  ÚretÚarg_dictÚarg_idxÚiÚarg_names	            r"   rm  zCompiledKernel.launch_metadataõ  sÊ   € Ü�=‰=×*Ñ*Ð2ØÜ §	¡	°t·}±}ÐPVÑWÓXˆÜ˜$Ÿ(™(¤IÔ.°$·(±(·+±+×2MÑ2MÐ2UØˆJØˆØˆÜ$ T§X¡X§[¡[×%:Ñ%:Ó;ò 	‰KˆAˆxØ!% g¡ˆH�XÑØ�q‰L‰Gð	ð 	�‰�—‘—‘×+Ñ+¨d°D·M±MÀ8Ð-LÔMØˆ
r#   c                ó<   ‡ ‡— ‰ j                  «        d dœˆˆ fd„
}|S )N)rj  c                ó|  •— | €=t         j                  j                  «       }t         j                  j                  |«      }  ‰j                  ‰| g|¢­Ž } ‰j
                  ‰d   ‰d   ‰d   | ‰j                  ‰j                  |t        j                  j                  t        j                  j                  g	|¢­Ž  y )Nr   r   r   )r   r©   rZ  Úget_current_streamrm  r\  rT  rP  r   rk  rl  Úlaunch_exit_hook)rj  r8  r¬   rm  rn  r>   s       €€r"   Úrunnerz*CompiledKernel.__getitem__.<locals>.runner  s�   ø€ Øˆ~ÜŸ™×9Ñ9Ó;�ÜŸ™×9Ñ9¸&ÓA�Ø2˜d×2Ñ2°4¸ÐGÀ$ÒGˆOØˆD�H‰H�T˜!‘W˜d 1™g t¨A¡w°¸¿¹Àt×G[ÑG[Ð]lÜ—]‘]×4Ñ4´e·m±m×6TÑ6Tð]ØW[ô]r#   )rd  )r>   rn  rx  s   `` r"   Ú__getitem__zCompiledKernel.__getitem__  s   ù€ Ø×ÑÔà!%÷ 	]ð ˆr#   )	r-   rf   rg   rB   rd  rg  rm  ry  Ú__classcell__)rh  s   @r"   r  r  »  s   ø„ òò:eô,.ò
ör#   r  )rÊ   ÚBaseExceptionra   )rò   r
   rc   r	   )3Ú
__future__r   rL   rõ   Ú_C.libtritonr   r   r   Úbackends.compilerr   r	   r
   r   r   r   Úruntime.autotunerr   Úruntime.cacher   r   r   Úruntime.driverr   Útools.disasmr   Úpathlibr   r   Ú	functoolsr�   rš   rÓ   Úptx_prototype_patternrr   Úptx_arg_type_patternru   r   r%   rj   Ú	lru_cacher¦   r¨   r´   rÂ   rÑ   r(  rú   r0  r0   r<  r  rh   r#   r"   ú<module>rˆ     s  ðÝ "Û Û ß >Ý Ý (ß 6ß !Ý .ß UÑ UÝ #Ý #Ý Û 	Û Û 	Û Û ð YÐ à	Ð ðÐ ð ,Ð à	ÐðÐ ò

÷#ñ #÷L'ñ 'ðT €×ÑÓñ 1ó ð 1ðF €×ÑÓñOó ðOò,ó#$÷L$
ñ $
óNN5ób÷)ñ )ô 
ˆdô 
÷Sò Sr#   