Ë
    6^(h&H  ã                  ó’  — 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 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m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jJ                  «       d„ «       Z&d„ Z'd"d„Z(d#d„Z)d„ Z* G d„ d«      Z+ G d„ de,«      Z- G d „ d!«      Z.y)$é    )ÚannotationsNé   )Úget_cache_invalidating_env_varsÚir)Úbackends)Ú	GPUTarget)Ú__version__)ÚOutOfResources)Úget_cache_managerÚget_dump_managerÚget_override_manager)Údriver)Úget_sassé   )Úast_to_ttir)Ú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 {[^}]+}Ú Ú*r   )ÚreÚsearchÚsubÚconvert_type_reprÚgroup)ÚxÚmatchÚtmas      úV/var/www/skyplay_api_hub/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                ó²  — || _         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)ÚfnÚextÚ__name__ÚnameÚ	signatureÚdictÚ	constantsÚitemsÚ
isinstanceÚstrÚ	arg_namesÚindexÚtupleÚattrsÚ	enumerateÚsplitÚstripÚkeysÚ	TypeError)Úselfr'   r+   Ú
constexprsr4   Ú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   Ã3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)ÚhasattrrA   r0   )r   s    r    ú<lambda>z ASTSource.hash.<locals>.<lambda>K   s   € ¬7°1°kÔ+B˜AŸK™K€ ÌÈAË€ r!   ú-úutf-8)Úsortedr+   r.   Újoinr-   r'   rA   r0   r4   ÚhashlibÚsha256ÚencodeÚ	hexdigest)r:   r<   r=   Ú
sorted_sigÚget_keyÚconstants_keyÚkeys          r    ÚhashzASTSource.hashI   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                ó8   — t        | j                  | ||||¬«      S )N)ÚcontextÚoptionsÚcodegen_fnsÚ
module_map)r   r'   ©r:   rS   rT   rU   rR   s        r    Úmake_irzASTSource.make_irP   s!   € Ü˜4Ÿ7™7 D°'À7ÐXcØ&0ô2ð 	2r!   c                ó   — t        «       S ©N)r,   ©r:   s    r    Úparse_optionszASTSource.parse_optionsT   s	   € Ü‹vˆr!   ©NN©ÚreturnÚNone©r)   Ú
__module__Ú__qualname__r>   rP   rW   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  | _        |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(                  j+                  «       }	d|	z   | _        | j(                  j-                  |	«      }
| j(                  j/                  |
«      }t!        |«      D ��ci c]  \  }}||“Œ
 c}}| _        y c c}}w c c}}w )Nr   r   r   ú@)Úpathr   Úsuffixr(   Ú	read_textÚsrcr   Úload_dialectsr   r   Úprototype_patternÚ	MULTILINEr   r*   ÚfindallÚarg_type_patternr5   r   r+   Úparse_mlir_moduleÚmoduleÚget_entry_func_nameÚget_functionÚget_function_signature)r:   rh   rR   Úbackendr   r+   Útypesr<   ÚtyÚfn_nameÚfuncOpÚfunc_tys               r    r>   zIRSource.__init__Z   sc  € ØˆŒ	Ü�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 )NrE   )rH   rI   rk   rJ   rK   rZ   s    r    rP   zIRSource.hashr   s'   € Ü�~‰~˜dŸh™hŸo™o¨gÓ6Ó7×AÑAÓCÐCr!   c                ó<   — || j                   _        | j                   S rY   )rr   rR   rV   s        r    rW   zIRSource.make_iru   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(   rr   Úget_int_attrr,   )r:   r€   s     r    r[   zIRSource.parse_optionsy   sF   € Ø�8‰8�wÒØŸ™×0Ñ0°ÓAˆIØÐ(ÐSÐ*SÓSÐ(Ø Ð+Ð+Ü‹vˆr!   Nr`   rc   r!   r    re   re   X   s   „ òEò0Dòór!   re   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   Úlanguageztriton.language.rD   )ÚpkgutilÚosrh   ÚdirnameÚabspathÚ__file__ÚopenrH   rI   ÚreadrK   rG   Úwalk_packagesÚmodule_finderÚ	find_specr*   ÚoriginÚ	sysconfigÚget_config_varr6   ÚupdateÚappendr	   )r‹   ÚTRITON_PATHÚcontentsÚfÚpath_prefixesrh   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                óì   — |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   rq   rR   r   rj   Ú
read_bytes)Ú	full_namer(   rR   rr   s       r    Úparserª   ¥   sx   € Ø
ˆf‚}˜˜wšÜ×%Ñ% i°Ó9ˆØ ˆŒØˆØ
ˆf‚}˜˜uš¨¨xªÜ�I‹×(Ñ(Ó*Ð*Ø
ˆg‚~˜ šÜ�I‹×)Ñ)Ó+Ð+ð (r!   c                ó*  ‡— t        j                  dd«      dk(  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!"
    ÚTRITON_FRONT_END_DEBUGGINGÚ0Ú1Nz"/triton/compiler/code_generator.pyz/ast.pyú/c              3  ó‚   •K  — | ]6  }‰j                   j                  j                  j                  |«      sŒ3|–— Œ8 y ­wrY   )Ú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Œ   ÚgetenvÚ	__cause__Úfilter_tracebackÚ__context__ÚreplaceÚsepÚ__traceback__Úanyr™   Útb_nextÚzip)ÚeÚ	BAD_FILESÚbad_fileÚframesÚ	cur_frameÚ
next_framer¶   s         @r    rº   rº   °   s  ø€ ô 
‡y�yÐ-¨sÓ3°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                óò	  — |€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&                  |j)                  d«      «      j+                  «       }	t-        |	«      }
t.        j0                  j3                  dd«      dk(  }t.        j0                  j3                  dd«      dk(  }t.        j0                  j3                  d	d«      dk(  }|rt5        | j                  «       «      nd }|rt7        | j                  «       «      nd }| j8                  d d
 }|› d�}|
j;                  |«      xs i }|j3                  |«      }t.        j0                  j3                  dd«      dk(  }|s|�t=        | ||	«      S |	|dœ|j>                  ¥|¥}t@        |d<   t        «       }|jC                  ||«       tE        |jG                  «       «      jI                  | jJ                  «      }|r|dz  }t        | t        «      s:t        j                  «       }t        jL                  |«       |jM                  |«       |jO                  |«      }|jQ                  «       }	 | jS                  |||«      }t.        j0                  j3                  dd «      }tE        |j#                  «       «      |d  D ]«  \  }} |||«      }|› d|› �} |�.|jY                  | «      x}!�t[        d|!› �«       t]        |!||«      }|r|dv r|
j_                  || «      || <   |�|j_                  || «       ||k(  r0|
jY                  | «      }"|ja                  |"«       t[        d|"› �«       |}Œ­ |
j_                  tc        jd                  |tf        ¬«      |d¬«      ||<   |
ji                  ||«       t.        j0                  j3                  dd«      dk(  s|jk                  «        t=        | ||	«      S # tT        $ r}tW        |«       ‚ d }~ww xY w)Nz target must be of GPUTarget typez'source must be either AST or a filepathrD   rE   ÚTRITON_KERNEL_OVERRIDEr­   r®   ÚTRITON_KERNEL_DUMPÚTRITON_STORE_BINARY_ONLYé–   ú.jsonÚTRITON_ALWAYS_COMPILE)rP   ÚtargetÚtriton_versionr   Ú
USE_IR_LOCr‡   z
Overriding kernel with file )r¦   r§   ÚjsonzCreating new locations for )ÚdefaultF)ÚbinaryÚTRITON_ENABLE_ASAN)6r   ÚactiveÚget_current_targetr/   r   Úmake_backendr#   r0   r   rR   re   r[   r,   r   r¢   rP   rF   r.   rH   rI   rJ   rK   r   rŒ   ÚenvironÚgetr   r   r*   Ú	get_groupÚCompiledKernelÚ__dict__r	   Ú
add_stagesÚlistr8   r2   r(   rl   Úget_codegen_implementationÚget_module_maprW   Ú	Exceptionrº   Úget_fileÚprintrª   ÚputÚcreate_location_snapshotrÒ   ÚdumpsÚvarsÚ	put_groupÚdisable_multithreading)#rk   rÏ   rS   rv   Ú	ir_sourcerR   Úextra_optionsÚenv_varsrO   rP   Úfn_cache_managerÚenable_overrideÚenable_ir_dumpÚstore_only_binaryÚfn_override_managerÚfn_dump_managerÚ	file_nameÚmetadata_filenameÚmetadata_groupÚmetadata_pathÚalways_compileÚmetadataÚstagesÚfirst_stagerT   rU   rr   rÂ   Ú
use_ir_locr(   Ú
compile_irÚnext_moduleÚir_filenamer©   Úir_full_names#                                      r    Úcompiler  Ö   sš  € Ø€~Ü—‘×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Ü(¨Ó.Ðô —j‘j—n‘nÐ%=¸sÓCÀsÑJ€OÜ—Z‘Z—^‘^Ð$8¸#Ó>À#ÑE€NÜŸ
™
Ÿ™Ð'AÀ3ÓGÈ3ÑNÐÙ>MÔ.¨s¯x©x«zÔ:ÐSWÐÙ6DÔ& s§x¡x£zÔ2È$€Oð
 —‘˜˜#�€IØ$˜+ UÐ+ÐØ%×/Ñ/Ð0AÓBÒHÀb€NØ"×&Ñ&Ð'8Ó9€MÜ—Z‘Z—^‘^Ð$;¸SÓAÀSÑH€NÙ˜mÐ7ä˜c >°4Ó8Ð8ð Øñð ×
Ñ
ðð ð	€Hô "-€HÐÑä‹V€FØ×Ñ�v˜wÔ'Ü�v—{‘{“}Ó%×+Ñ+¨C¯G©GÓ4€KáØ�qÑˆô �cœ8Ô$Ü—*‘*“,ˆÜ
×Ñ˜Ô!Ø×Ñ˜gÔ&à×4Ñ4°WÓ=€KØ×'Ñ'Ó)€JðØ—‘˜W k°:¸wÓGˆô —‘—‘ ¨dÓ3€JÜ §¡£Ó/°°Ð=ò ‰ˆˆZÙ  ¨Ó2ˆØ"˜ 1 S EÐ*ˆØÐ+Ð>Q×>ZÑ>ZÐ[fÓ>gÐ1g°Ð0tÜÐ2°9°+Ð>Ô?Ü 	¨3°Ó8ˆKá! sÐ.HÑ'HØ*:×*>Ñ*>¸{ÈKÓ*XˆN˜;Ñ'ØÐ&Ø×Ñ ¨[Ô9à˜ÒØ+×4Ñ4°[ÓAˆLØ×0Ñ0°Ô>ÜÐ/°¨~Ð>Ô?Ø‰ð!ð$ )9×(<Ñ(<¼T¿Z¹ZÈÔZ^Ô=_ÐarØDIð )=ó )K€NÐ$Ñ%à×ÑÐ0°.ÔAô �:‰:�>‰>Ð.°Ó4¸Ò;Ø×&Ñ&Ô(ä˜#˜~¨tÓ4Ð4øôK ò Ü˜ÔØûðús   ÍS Ó	S6Ó%S1Ó1S6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ÚRuntimeErrorrv   )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d„Zd„ Zy)ÚLazyDictc                ó    — || _         g | _        y rY   )ÚdataÚextras)r:   r  s     r    r>   zLazyDict.__init__I  s   € ØˆŒ	Øˆ�r!   c                ó¤   — | j                   D ]  \  }}| j                   ||Ž z  | _        Œ | j                   j                  «        | j                  S rY   )r  r  Úclear©r:   ÚfuncÚargss      r    rÚ   zLazyDict.getM  sG   € ØŸ+™+ò 	0‰JˆD�$ØŸ	™	¡D¨$ KÑ/ˆD�Ið	0à�‰×ÑÔØ�y‰yÐr!   c                ó>   — | j                   j                  ||f«       y rY   )r  r™   r  s      r    ÚaddzLazyDict.addS  s   € Ø�‰×Ñ˜D $˜<Õ(r!   Nr]   )r)   ra   rb   r>   rÚ   r  rc   r!   r    r	  r	  G  s   „ òóó)r!   r	  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:   rO   Úvalues      r    Ú__missing__zAsmDict.__missing__Y  s6   € à�&Š=Ü˜T '™]Ó+‰EäÐ.°Ñ4Ó5Ð5àˆˆS‰	Øˆr!   N)r)   ra   rb   r  rc   r!   r    r  r  W  s   „ ór!   r  c                  ó>   ‡ — e Zd ZdZdZd„ Zd„ Zˆ fd„Zd„ Zd„ Z	ˆ xZ
S )rÜ   Nc           	     óÔ  — 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>m  s$   è ø€ Ò`©$¨!¨QÈAÏJÉJÐW^ÕL_œd 1ŸgÑ`ùs   ‚,œ,Úcluster_dimsrÏ   rv   ÚarchÚ	warp_sizeÚKernelMetadatarÍ   r   rc   )Úcollectionsr  Únextr.   rÒ   Úloadsrj   r3   r   rF   rß   r8   rù   rØ   rÏ   Úpack_metadataÚpacked_metadatark   rP   r*   r´   r   Ú
binary_extr  ri   r¨   ÚasmÚkernelrr   Úfunction)r:   rk   rö   rP   r  r÷   rù   rÏ   r$  rv   r  r   Ú	asm_filesr*  Úfiles                  r    r>   zCompiledKernel.__init__k  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                óB  — | j                   �y t        j                  j                  «       }t        j                  j	                  | j
                  | j                  «      | _        t        j                  j                  j                  |«      d   }| 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                  |«      \  | _         | _        | _        | _        y )NÚmax_shared_memzshared memoryÚ	tmem_sizei   ztensor memory)rr   r   rÖ   Úget_current_deviceÚlauncher_clsrk   rù   ÚrunÚutilsÚget_device_propertiesÚsharedr
   rB   r2  Úload_binaryr*   r,  r-  Ún_regsÚn_spills)r:   ÚdeviceÚ
max_sharedÚmax_tmem_sizes       r    Ú_init_handleszCompiledKernel._init_handlesˆ  s!  € Ø�;‰;Ð"ØÜ—‘×1Ñ1Ó3ˆä—=‘=×-Ñ-¨d¯h©h¸¿¹ÓFˆŒä—]‘]×(Ñ(×>Ñ>¸vÓFÐGWÑXˆ
Ø�=‰=×Ñ *Ò,Ü  §¡×!5Ñ!5°zÀ?ÓSÐSÜ�4—=‘= +Ô.°4·=±=×3JÑ3JÐ3VàˆMØ�}‰}×&Ñ&¨Ò6Ü$ T§]¡]×%<Ñ%<¸mÈ_Ó]Ð]äAGÇÁ×ATÑAT×A`ÑA`Ø�I‰I�t—{‘{ D§M¡M×$8Ñ$8¸&óBBÑ>ˆŒ�T”] D¤K°µr!   c                óL   •— |dk(  r| j                  «        t        ‰| �	  |«      S )Nr5  )r?  ÚsuperÚ__getattribute__)r:   r*   Ú	__class__s     €r    rB  zCompiledKernel.__getattribute__›  s&   ø€ Ø�5Š=Ø×ÑÔ Ü‰wÑ'¨Ó-Ð-r!   c                óä  — t         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*   r-  Ústreamr   r   )rÜ   Úlaunch_enter_hookr	  r*   r-  r/   rk   r#   r'   Úlaunch_metadatar5   r1   r  rù   )	r:   ÚgridrE  r  ÚretÚarg_dictÚarg_idxÚiÚarg_names	            r    rG  zCompiledKernel.launch_metadata   sÆ   € Ü×+Ñ+Ð3ØÜ §	¡	°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)rE  c                óT  •— | €=t         j                  j                  «       }t         j                  j                  |«      }  ‰j                  ‰| g|¢­Ž } ‰j
                  ‰d   ‰d   ‰d   | ‰j                  ‰j                  |t        j                  t        j                  g	|¢­Ž  y )Nr   r   r   )r   rÖ   r3  Úget_current_streamrG  r5  r-  r)  rÜ   rF  Úlaunch_exit_hook)rE  r  r<  rG  rH  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Ü#×5Ñ5´~×7VÑ7Vð_ØY]ô_r!   )r?  )r:   rH  rR  s   `` r    Ú__getitem__zCompiledKernel.__getitem__®  s   ù€ Ø×ÑÔà!%÷ 	_ð ˆr!   )r)   ra   rb   rF  rQ  r>   r?  rB  rG  rS  Ú__classcell__)rC  s   @r    rÜ   rÜ   d  s+   ø„ ð ÐØÐòò:Bô&.ò
ör!   rÜ   )rÂ   ÚBaseExceptionr\   )/Ú
__future__r   rH   rÒ   Ú_C.libtritonr   r   r   Úbackends.compilerr   r   r	   Úruntime.autotunerr
   Úruntime.cacher   r   r   Úruntime.driverr   Útools.disasmr   Úcode_generatorr   Úpathlibr   r   Ú	functoolsrŒ   r–   Úptx_prototype_patternrm   Úptx_arg_type_patternrp   r   r#   re   Ú	lru_cacher¢   rª   rº   r  rØ   r	  r,   r  rÜ   rc   r!   r    ú<module>rc     sÜ   ðÝ "Û Û ß >Ý Ý )Ý Ý .ß UÑ UÝ #Ý #å 'Ý Û 	Û Û 	Û ð YÐ à	Ð ðÐ ð ,Ð à	ÐðÐ ò

÷!ñ !÷H&ñ &ðR €×ÑÓñ 1ó ð 1òF,ó#$óLf5òR÷)ñ )ô 
ˆdô 
÷Uò Ur!   