
    'Th&H                       S SK Jr  S SKrS SKrSSKJrJr  SSKJr  SSKJ	r	  SSK
Jr  SSKJr  SS	KJrJrJr  SS
KJr  SSKJr  SSKJr  S SKJr  S SKrS SKrS SKrS SKrSrS\0rSr S\ 0r!S r" " S S5      r# " S S5      r$\RJ                  " 5       S 5       r&S r'S"S jr(S#S jr)S r* " S S5      r+ " S S\,5      r- " S  S!5      r.g)$    )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                    [         R                  " SU 5      n[         R                  " SU 5      nUb  g[         R                  " SSU 5      n Ub  S[        UR	                  S5      5      -   $ U $ )Nz!tt\.ptr<([^,]+)ztt.nv_tma_desc = 1	nvTmaDescz {[^}]+} *r   )researchsubconvert_type_reprgroup)xmatchtmas      P/var/www/auris/envauris/lib/python3.13/site-packages/triton/compiler/compiler.pyr   r   '   sd     II)1-E
)))1
-C

{B"A&u{{1~666H    c                  4    \ rS rSrSS	S jjrS rS rS rSrg)
	ASTSource4   Nc                   Xl         SU l        UR                  U l        X l        [        5       U l        Ubq  UR                  5        H]  u  pV[        U[        5      (       a  UR                  R                  U5      4OUn[        U[        5      (       d   eX`R                  U'   M_     U=(       d
    [        5       U l        [        U R                  [        5      (       aM  [        U R                  R                  S5      5       VVs0 s H  u  pVXVR!                  5       _M     snnU l        g U R                  R#                  5        H#  n[        U[        5      (       a  M  [%        S5      e   g s  snnf )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,   
constexprsr5   kvs          r    __init__ASTSource.__init__6   s   KK	"!"((*1;As1C1CR\\''*-!!U++++$%q! + _df
dnnc**7@AUAUVYAZ7[\7[tqal7[\DN^^((*!!S))#$CDD + ]s   E.c           	        [        U R                  R                  5       5       VVs/ s H  u  pUPM	     nnnS nSR                  [        U R                  R                  5       5       VVs/ s H  u  pU" U5      PM     snn5      nU R
                  R                   S[        U R                  5       SU SU 3n[        R                  " UR                  S5      5      R                  5       $ s  snnf s  snnf )Nc                R    [        U S5      (       a  U R                  $ [        U 5      $ )N	cache_key)hasattrrC   r1   )r   s    r    <lambda> ASTSource.hash.<locals>.<lambda>K   s     71k+B+BAKKNANr!   -utf-8)sortedr,   r/   joinr.   r(   rC   r1   r5   hashlibsha256encode	hexdigest)r;   r=   r>   
sorted_sigget_keyconstants_keykeys          r    hashASTSource.hashI   s    $*4>>+?+?+A$BC$BDAa$B
CN@T@T@V9W!X9W'!*9W!XY""#1S_$5Qzl!M?S~~cjj12<<>>	 D!Xs   C/,C5
c           	     .    [        U R                  XXUS9$ )N)contextoptionscodegen_fns
module_map)r   r(   r;   rW   rX   rY   rV   s        r    make_irASTSource.make_irP   s    477D7&02 	2r!   c                    [        5       $ N)r-   r;   s    r    parse_optionsASTSource.parse_optionsT   s	    vr!   )r5   r.   r)   r(   r+   r,   NNreturnNone	r*   
__module____qualname____firstlineno__r?   rS   r[   r`   __static_attributes__ r!   r    r#   r#   4   s    E&?2r!   r#   c                  ,    \ rS rSrS rS rS rS rSrg)IRSourceX   c                   Xl         [        U5      nUR                  SS  U l        UR	                  5       U l        [        R                  " U5        UR                  U5        U R                  S:X  a  [        R                  " [        U R                     U R
                  [        R                  5      nUR                  S5      U l        UR                  S5      n[        R                  " [        U R                     U5      n[!        U5       VVs0 s H  u  pxU[#        U5      _M     snnU l        g [        R&                  " U R                   U5      U l        U R(                  R+                  5       n	SU	-   U l        U R(                  R-                  U	5      n
U R(                  R/                  U
5      n[!        U5       VVs0 s H  u  pxXx_M	     snnU l        g s  snnf s  snnf )Nr   r   r   @)pathr   suffixr)   	read_textsrcr   load_dialectsr   r   prototype_pattern	MULTILINEr   r+   findallarg_type_patternr6   r   r,   parse_mlir_modulemoduleget_entry_func_nameget_functionget_function_signature)r;   rq   rV   backendr   r,   typesr=   tyfn_namefuncOpfunc_tys               r    r?   IRSource.__init__Z   sg   	Dz;;qr?>>#
!g& 88uII/9488R\\READIAIJJ/99EEDMeDTUDT51a!22!66DTUDN..tyy'BDKkk557GgDI[[--g6Fkk88@G1:71CD1Cae1CDDN V Es   G Gc                |    [         R                  " U R                  R                  S5      5      R	                  5       $ )NrH   )rK   rL   rt   rM   rN   r_   s    r    rS   IRSource.hashr   s'    ~~dhhoog67AACCr!   c                :    X@R                   l        U R                   $ r^   )r{   rV   rZ   s        r    r[   IRSource.make_iru   s    %{{r!   c                    U R                   S:X  a)  U R                  R                  S5      nUc   S5       eSU0$ [        5       $ )Nttgirzttg.num-warpsz'Unable to parse ttg.num-warps attribute	num_warps)r)   r{   get_int_attrr-   )r;   r   s     r    r`   IRSource.parse_optionsy   sF    88w00AI(S*SS(++vr!   )r)   r{   r+   rq   r,   rt   Nrf   rk   r!   r    rm   rm   X   s    E0Dr!   rm   c                    SS K n [        R                  R                  [        R                  R                  [        R                  R	                  [
        5      5      5      n/ n[        [
        S5       nU[        R                  " UR                  5       5      R                  5       /-  nS S S 5        [        R                  R                  US5      S4[        R                  R                  US5      S4/nU H  u  pVU R                  U/US9 H{  n[        UR                  R                  UR                  5      R                   S5       nU[        R                  " UR                  5       5      R                  5       /-  nS S S 5        M}     M     [        R                  " 5       n["        R$                  " S5      R'                  S	5      S
   n	[        [        R                  R                  USSU	 35      S5       n UR                  S5      n
U
(       d  OUR)                  U
5        M,  S S S 5        UR+                  UR                  5       5        [        R                  R                  US5      nU R                  U/SS9 H{  n[        UR                  R                  UR                  5      R                   S5       nU[        R                  " UR                  5       5      R                  5       /-  nS S S 5        M}     [,         SR                  U5      -   $ ! , (       d  f       GNw= f! , (       d  f       GM.  = f! , (       d  f       GN= f! , (       d  f       M  = f)Nr   rbcompilerztriton.compiler.r   ztriton.backends.)prefix
EXT_SUFFIX._Cz
libtriton.i   languageztriton.language.rG   )pkgutilosrq   dirnameabspath__file__openrK   rL   readrN   rJ   walk_packagesmodule_finder	find_specr+   origin	sysconfigget_config_varr7   updateappendr	   )r   TRITON_PATHcontentsfpath_prefixesrq   r   liblibtriton_hashr)   chunklanguage_paths               r    
triton_keyr      s   ''//"''//"''//(2K"LMKH	h	W^^AFFH-779:: 
 
k:	.0BC	k:	.0BCM &(($(?Cc''11#((;BBDIQW^^AFFH5??ABB JI @ & ^^%N

"
"<
0
6
6s
;B
?C	bggll;
3%.@A4	HAFF7OE!!%(	  
I OON,,./GGLLj9M$$m_=O$P###--chh7>>E1;;=>>H FE Q ]chhx0007 
	 JI 
I	H FEs0   47L7L-=.M ?7M
L*-
L= 
M
M!	c                    US:X  d  US:X  a  [         R                  " X5      nX#l        U$ US:X  d  US:X  d  US:X  a  [        U 5      R	                  5       $ US:X  d  US:X  a  [        U 5      R                  5       $ g )Nr&   r   llirr   amdgcncubinhsaco)r   rz   rV   r   rs   
read_bytes)	full_namer)   rV   r{   s       r    parser      st    
f}w%%i9 
f}uxI((**
g~I))++ (r!   c                P  ^ [         R                  " SS5      S:X  a  gU R                  b  [        U R                  5        U R                  b  [        U R                  5        SS/nU Vs/ s H"  o"R                  S[         R                  5      PM$     nnU R                  m/ nTb<  [        U4S jU 5       5      (       d  UR                  T5        TR                  mTb  M<  [        X3S	S 5       H  u  pEXTl
        M     U(       d  SU l        gSUS
   l
        US   U l        gs  snf )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     >#    U  H=  nTR                   R                  R                  R                  U5      (       d  M9  Uv   M?     g 7fr^   )tb_framef_codeco_filenameendswith).0r   tbs     r    	<genexpr>#filter_traceback.<locals>.<genexpr>   s1     Vi2;;+=+=+I+I+R+RST+U11is
   8A	A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    
yy-s3s:{{%}} ' 	-I @IIy8!!#rvv.yII	
BF
.ViVVVMM"ZZ .
 $'vabz#:& $; !r
 )! Js   *)D#c                |
   Uc  [         R                  R                  5       n[        U[        5      (       d   S5       e[        U5      n[        U [        5      (       + nU(       a=  [        U [        5      (       d   S5       e[        R                  " 5       n[        XU5      n U R                  5       nUR                  [        U=(       d
    [        5       40 UD65      n[        5       n[        5        SU R                  5        SUR                  5        SUR                  5        S[        [!        UR#                  5       5      5       3	n[$        R&                  " UR)                  S5      5      R+                  5       n	[-        U	5      n
[.        R0                  R3                  SS5      S:H  n[.        R0                  R3                  SS5      S:H  n[.        R0                  R3                  S	S5      S:H  nU(       a  [5        U R                  5       5      OS nU(       a  [7        U R                  5       5      OS nU R8                  S S
 nU S3nU
R;                  U5      =(       d    0 nUR3                  U5      n[.        R0                  R3                  SS5      S:H  nU(       d  Ub  [=        U UU	5      $ U	US.UR>                  EUEn[@        US'   [        5       nURC                  UU5        [E        URG                  5       5      RI                  U RJ                  5      nU(       a  US-  n[        U [        5      (       d<  [        R                  " 5       n[        RL                  " U5        URM                  U5        URO                  U5      nURQ                  5       n U RS                  UUUW5      n[.        R0                  R3                  SS 5      n[E        UR#                  5       5      US   H  u  nnU" UU5      nU SU 3n Ub/  URY                  U 5      =n!b  [[        SU! 35        []        U!UU5      nU(       a  US;   a  U
R_                  UU 5      UU '   Ub  UR_                  UU 5        UU:X  a0  U
RY                  U 5      n"URa                  U"5        [[        SU" 35        UnM     U
R_                  [b        Rd                  " U[f        S9USS9UU'   U
Ri                  UU5        [.        R0                  R3                  SS5      S:X  d  URk                  5         [=        U UU	5      $ ! [T         a  n[W        U5        e S nAff = f)Nz target must be of GPUTarget typez'source must be either AST or a filepathrG   rH   TRITON_KERNEL_OVERRIDEr   r   TRITON_KERNEL_DUMPTRITON_STORE_BINARY_ONLY   .jsonTRITON_ALWAYS_COMPILE)rS   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_targetr0   r   make_backendr#   r1   r   rV   rm   r`   r-   r   r   rS   rI   r/   rK   rL   rM   rN   r   r   environgetr   r   r+   	get_groupCompiledKernel__dict__r	   
add_stageslistr9   r3   r)   ru   get_codegen_implementationget_module_mapr[   	Exceptionr   get_fileprintr   putcreate_location_snapshotr   dumpsvars	put_groupdisable_multithreading)#rt   r   rW   r   	ir_sourcerV   extra_optionsenv_varsrR   rS   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_stagerX   rY   r{   r   
use_ir_locr)   
compile_irnext_moduleir_filenamer   ir_full_names#                                      r    compiler     s   ~113fi((L*LL(6"GsI..I#s##N%NN#**,sW-%%'M##D):DF$Lm$LMG.0H\N!CHHJ<q(8',,.9I3vV^VdVdVfOgKhJi
jC>>#**W-.88:D(. jjnn%=sCsJOZZ^^$8#>#EN

'A3G3N>M.sxxz:SW6D&sxxz2$O
 #I$+U+%//0ABHbN"&&'89MZZ^^$;SASHNm7c>488  

 	H "-HVFvw'v{{}%++CGG4Kq c8$$**,
!g&44W=K'')JWk:wG d3J/=Z 2"1SE*+>Q>Z>Z[f>g1g0t29+>?	38K!s.H'H*:*>*>{K*XN;'&[9+44[AL00>/~>?! >$ )9(<(<TZZZ^=_arDI )= )KN$%0.A ::>>.4;&&(#~t44K  s   T   
T;*T66T;c                .   [         R                  " 5        Vs/ s H0  oR                  R                  U 5      (       d  M$  UR                  PM2     nn[	        U5      S:w  a'  [        [	        U5       SU R                   SU S35      eUS   " U 5      $ s  snf )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[#4a

8R8RSY8Zzqzz#4G[
7|q7|n=fnn=MSQXPYYuvx 	x1:f	 \s
   #BBc                  *    \ rS rSrS rSS jrS rSrg)LazyDictiG  c                    Xl         / U l        g r^   dataextras)r;   r  s     r    r?   LazyDict.__init__I  s    	r!   c                    U R                    H  u  pU R                  U" U6 -  U l        M     U R                   R                  5         U R                  $ r^   )r  r  clearr;   funcargss      r    r   LazyDict.getM  s@    ++JD		D$K/DI &yyr!   c                <    U R                   R                  X45        g r^   )r  r   r  s      r    addLazyDict.addS  s    D<(r!   r  Nrc   )r*   rg   rh   ri   r?   r   r$  rj   rk   r!   r    r  r  G  s    )r!   r  c                      \ rS rSrS rSrg)AsmDictiW  c                T    US:X  a  [        U S   5      nO[        SU-  5      eX U'   U$ )Nsassr   zUnknown key: '%s')r   KeyError)r;   rR   values      r    __missing__AsmDict.__missing__Y  s4    &=T']+E.455S	r!   rk   N)r*   rg   rh   ri   r,  rj   rk   r!   r    r'  r'  W  s    r!   r'  c                  H   ^  \ rS rSrSrSrS rS rU 4S jrS r	S r
SrU =r$ )	r   id  Nc           	        SSK Jn  [        S UR                  5        5       5      n[        R
                  " UR                  5       5      n[        US   5      US'   US   n[        US   US   US   5      US'   U" S	[        [        UR                  5       5      5      5      nU" S0 UD6U l        [        U R                  R                  5      n	U	R                  U R                  5      U l        Xl        X0l        U R                  R&                  U l        UR                  5        V
Vs/ s H(  u  pU
R)                  S
5      (       a  M  [+        U5      PM*     nn
nU	R,                  n[/        U Vs0 s HD  nUR0                  SS  UR0                  SS  U:X  a  UR3                  5       OUR                  5       _MF     sn5      U l        U R4                  U   U l        S U l        S U l        g s  snn
f s  snf )Nr   )
namedtuplec              3  l   #    U  H*  u  pUR                  S 5      (       d  M  [        U5      v   M,     g7f)r   N)r   r   )r   cps      r    r   *CompiledKernel.__init__.<locals>.<genexpr>m  s(     `2H$!AJJW^L_gd1gg2Hs   44cluster_dimsr   r   arch	warp_sizeKernelMetadatar   r   rk   )collectionsr0  nextr/   r   loadsrs   r4   r   rI   r   r9   r  r   r   pack_metadatapacked_metadatart   rS   r+   r   r   
binary_extr'  rr   r   asmkernelr{   function)r;   rt   r  rS   r0  r  r  r   r8  r   r2  r3  	asm_filesr>  files                  r    r?   CompiledKernel.__init__k  s   *`.2F2F2H`a::m5578#(.)A#B (#&vi'8&.&Q\J]^#$4fT(--/=R6ST&22t}}334&44T]]C	MM&&	)7)=)=)?[)?qzzRYGZWT!W)?	[''
!
! KKO$++ab/Z2OT__.UYUcUcUee!
  hhz*  \
s   )G*	G*0AG0c                T   U R                   b  g [        R                  R                  5       n[        R                  R	                  U R
                  U R                  5      U l        [        R                  R                  R                  U5      S   nU R                  R                  U:  a!  [        U R                  R                  US5      e[        U R                  S5      (       aT  U R                  R                  b=  SnU R                  R                  U:  a!  [        U R                  R                  US5      e[        R                  R                  R                  U R                  U R                   U R                  R                  U5      u  U l         U l        U l        U l        g )Nmax_shared_memzshared memory	tmem_sizei   ztensor memory)r{   r   r   get_current_devicelauncher_clsrt   r  runutilsget_device_propertiessharedr
   rD   rG  load_binaryr+   r@  rA  n_regsn_spills)r;   device
max_sharedmax_tmem_sizes       r    _init_handlesCompiledKernel._init_handles  s$   ;;"113==--dhhF]]((>>vFGWX
==*, !5!5z?SS4==+..4==3J3J3VM}}&&6$T]]%<%<m_]]AGATATA`A`IIt{{DMM$8$8&BB>T]DKr!   c                N   > US:X  a  U R                  5         [        TU ]	  U5      $ )NrJ  )rT  super__getattribute__)r;   r+   	__class__s     r    rX  CompiledKernel.__getattribute__  s&    5= w'--r!   c                   [         R                  c  g [        U R                  U R                  US.5      n[        U R                  [        5      (       a!  U R                  R                  R                  c  U$ 0 nSn[        U R                  R                  R                  5       H  u  pxX6   XX'   US-  nM     UR                  U R                  R                  R                  XR                  U45        U$ )N)r+   rA  streamr   r   )r   launch_enter_hookr  r+   rA  r0   rt   r#   r(   launch_metadatar6   r2   r$  r  )	r;   gridr\  r!  retarg_dictarg_idxiarg_names	            r    r^  CompiledKernel.launch_metadata  s    ++3		t}}PVWX$((I..$((++2M2M2UJ$TXX[[%:%:;KA!%HqLG < 	++dMM8-LM
r!   c                @   ^ ^ T R                  5         S S.UU 4S jjnU$ )N)r\  c                V  > U c=  [         R                  R                  5       n[         R                  R                  U5      n TR                  " TU /UQ76 nTR
                  " TS   TS   TS   U TR                  TR                  U[        R                  [        R                  /	UQ76   g )Nr   r   r   )r   r   rH  get_current_streamr^  rJ  rA  r=  r   r]  launch_exit_hook)r\  r!  rQ  r^  r_  r;   s       r    runner*CompiledKernel.__getitem__.<locals>.runner  s    ~99;99&A"224G$GOHHT!Wd1gtAwtG[G[]l#55~7V7V_Y]_r!   )rT  )r;   r_  rj  s   `` r    __getitem__CompiledKernel.__getitem__  s%    !% 	_ 	_ r!   )r?  rA  rS   r@  r  r{   rO  rP  r+   r=  rJ  rt   )r*   rg   rh   ri   r]  ri  r?   rT  rX  r^  rl  rj   __classcell__)rY  s   @r    r   r   d  s0     :B&.
 r!   r   )r   BaseExceptionrb   )/
__future__r   rK   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_patternrv   ptx_arg_type_patternry   r   r#   rm   	lru_cacher   r   r   r  r   r  r-   r'  r   rk   r!   r    <module>r}     s    "   >  )  . U U # # '  	  	  Y 	   , 	 

! !H& &R  1  1F,#$Lf5R) ) 
d 
U Ur!   