
    hT                       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
Jr  SSKJr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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&\RN                  " 5       S 5       r(\RN                  " 5       S 5       r)S r*S$S jr+ " S S5      r,S%S jr-S&S jr. " S S5      r/ " S  S!\05      r1 " S" S#5      r2g)'    )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                    [         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 {[^}]+} *   )researchsubconvert_type_reprgroup)xmatchtmas      R/var/www/fran/franai/venv/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         [        R                  U 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)fnr   TRITONlanguageext__name__name	signaturedict	constantsitems
isinstancestr	arg_namesindextupleattrs	enumeratesplitstripkeys	TypeError)selfr*   r0   
constexprsr9   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   F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)hasattrrG   r5   )r   s    r"   <lambda> ASTSource.hash.<locals>.<lambda>L   s     71k+B+BAKKNANr#   -utf-8)sortedr0   r3   joinr2   r*   rG   r5   r9   hashlibsha256encode	hexdigest)r?   rA   rB   
sorted_sigget_keyconstants_keykeys          r"   hashASTSource.hashJ   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           	     4    SSK Jn  U" U R                  XXUS9$ )Nr   )ast_to_ttir)contextoptionscodegen_fns
module_map)code_generatorrZ   r*   )r?   r\   r]   r^   r[   rZ   s         r"   make_irASTSource.make_irQ   s    /477D7&02 	2r#   c                    [        5       $ N)r1   r?   s    r"   parse_optionsASTSource.parse_optionsV   s	    vr#   )r9   r2   r-   r*   r,   r/   r0   NNreturnNone	r.   
__module____qualname____firstlineno__rC   rW   r`   re   __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)IRSourceZ   c                   Xl         [        U5      nUR                  SS  U l        [        R
                  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.                  R1                  5       n	SU	-   U l        U R.                  R3                  U	5      n
U R.                  R5                  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-   r   r+   r,   	read_textsrcr   load_dialectsr   r   prototype_pattern	MULTILINEr   r/   findallarg_type_patternr:   r   r0   parse_mlir_modulemoduleget_entry_func_nameget_functionget_function_signature)r?   rv   r[   backendr    r0   typesrA   tyfn_namefuncOpfunc_tys               r"   rC   IRSource.__init__\   sr   	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-G3c                |    [         R                  " U R                  R                  S5      5      R	                  5       $ )NrL   )rO   rP   ry   rQ   rR   rd   s    r"   rW   IRSource.hashu   s'    ~~dhhoog67AACCr#   c                :    X@R                   l        U R                   $ rc   )r   r[   )r?   r\   r]   r^   r[   s        r"   r`   IRSource.make_irx   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_attrr1   )r?   r   s     r"   re   IRSource.parse_options|   sF    88w00AI(S*SS(++vr#   )r-   r,   r   r/   rv   r0   ry   Nrk   rp   r#   r"   rr   rr   Z   s    E2Dr#   rr   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   r,   ztriton.language.rK   )pkgutilosrv   dirnameabspath__file__openrO   rP   readrR   rN   walk_packagesmodule_finder	find_specr/   origin	sysconfigget_config_varr;   updateappendr   )r   TRITON_PATHcontentsfpath_prefixesrv   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                Z    [         R                  R                  R                  U 5      S   $ )Nmax_shared_mem)r   activeutilsget_device_properties)devices    r"   r   r      s#    ==44V<=MNNr#   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   r   r[   r   rx   
read_bytes)	full_namer-   r[   r   s       r"   parser      st    
f}w%%i9 
f}uxI((**
g~I))++ (r#   c                X  ^ [         R                  R                  (       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!"
Nz"/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rc   )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   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Iy8!!#rvv.yII	
BF
.ViVVVMM"ZZ .
 $'vabz#:& $; !r
 )! Js   .)D'c                  <    \ rS rSrSS jrSS jrS	S jrS
S jrSrg)CompileTimer   c                b    [         R                   " 5       U l        S U l        / U l        S U l        g rc   )timestartir_initialization_endlowering_stage_endsstore_results_endrd   s    r"   rC   CompileTimer.__init__   s'     IIK
37"<> /3r#   c                8    [         R                   " 5       U l        g rc   )r   r   rd   s    r"   finished_ir_initialization'CompileTimer.finished_ir_initialization   s    %)YY["r#   c                d    U R                   R                  U[        R                  " 5       45        g rc   )r   r   r   )r?   
stage_names     r"   stage_finishedCompileTimer.stage_finished   s!      ''TYY[(ABr#   c                f   [         R                   " 5       nU R                  c  Xl        OXl        SS jn/ nU R                  nU R                   H  u  pVUR	                  XR" XF5      45        UnM!     [
        R                  " U" U R                  U R                  5      UU" X@R                  5      S9$ )Nc                *    Uc  g[        X-
  S-  5      $ )Nr   i@B )int)r   ends     r"   deltaCompileTimer.end.<locals>.delta   s    {w.//r#   )ir_initializationlowering_stagesstore_results)r   floatr   zfloat | Noneri   r   )r   r   r   r   r   r   CompileTimesr   )r?   	timestampr   lowering_stage_durationsstage_startr   	stage_ends          r"   r   CompileTimer.end   s    IIK	%%-)2&%."	0
 $& 00%)%=%=!J$++Z{9V,WX#K &> !!#DJJ0J0JK4-C-CD
 	
r#   )r   r   r   r   Nrh   )r   r5   ri   rj   )ri   zknobs.CompileTimes)	r.   rl   rm   rn   rC   r   r   r   ro   rp   r#   r"   r   r      s    41C
r#   r   c                :   [         R                  R                  nU(       a
  [        5       n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
R1                  S5      5      R3                  5       n[5        U5      n[         R                  R6                  n[         R                  R8                  n[         R                  R:                  nU(       a  [=        U R'                  5       5      OS nU(       a  [?        U R'                  5       5      OS nU R@                  S S nU S3nURC                  U5      =(       d    0 nURE                  U5      n[         R                  RF                  nU(       dI  UbF  [I        U UU5      nU(       a0  U" U URJ                  RM                  5       UWRO                  5       SS9  U$ UUS	.URP                  EU	En[R        US
'   [!        5       nURU                  UX RV                  5        [Y        UR[                  5       5      R]                  U R^                  5      nU(       a  US-  n[        U [        5      (       d<  [        R                  " 5       n[        R`                  " U5        URa                  U5        URc                  U5      nURe                  5       n U Rg                  UUUW5      nU(       a'  U SU R^                   3nURm                  UU5      UU'   OU S3nURm                  UU5      UU'   [         R                  Rn                  n U(       a:  U (       a3  URq                  U Rr                  5        [u        SU Rr                   35        U(       a  WRw                  5         [Y        UR+                  5       5      US   GH  u  n!n"U"" UU5      n#U SU! 3nUc@  URE                  SS 5      =n$(       a&  U$Ry                  SU! 35      (       a  [{        U$U!U5      n#O3UR}                  U5      =n%(       a  [u        SU% 35        [{        U%U!U5      n#U(       a  U!S;   a  URm                  U#U5      UU'   Ub  URm                  U#U5        U U!:X  a0  UR}                  U5      n&U#Rq                  U&5        [u        SU& 35        U#nU(       d  GM  WR                  U!5        GM     URm                  [        R                  " U[        S9USS9UU'   UR                  UU5        [         R                  R                  (       d  UR                  5         U(       a  U" U UUWRO                  5       SS9  [I        U UU5      $ ! [h         a  n[k        U5        e S nAff = f)Nz target must be of GPUTarget typez'source must be either AST or a filepathrK   rL      .jsonT)ry   metadatametadata_grouptimes	cache_hit)rW   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_targetr4   r
   make_backendr%   r5   r   r[   rr   re   r1   r   r   rW   rM   r3   rO   rP   rQ   rR   r   overridedump_irstore_binary_onlyr   r   r/   	get_groupgetalways_compileCompiledKernelr  _asdictr   __dict__r   
add_stagesr,   listr=   r7   r-   rz   get_codegen_implementationget_module_mapr`   	Exceptionr   put
use_ir_loccreate_location_snapshotrv   printr   r   r   get_filer   r	  dumpsvars	put_groupenable_asandisable_multithreading)'ry   r  r\   compilation_listenertimerr   	ir_sourcer[   extra_optionsenv_varsrV   rW   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_stager]   r^   r   r   ir_filenamer  r-   
compile_irnext_moduler  r   ir_full_names'                                          r"   compiler<    s    ,,55~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(. ''00O&&..N));;>M.sxxz:SW6D&sxxz2$O
 #I$+U+%//0ABHbN"&&'89M&&55Nm7S.$7 --/-iik 
  

 	H "-HVFvw5v{{}%++CGG4Kq c8$$**,
!g&44W=K'')JWk:wG
 "1SWWI.&6&:&:6;&O{#"7+&6&:&:6;&O{#""--JZ''1+CHH:67((*/=Z 2"1SE*&  (||M4@@@kFZFZ]^_b^c[dFeFe#Kg>-66{CCYC29+>?	38K!s.H'H*:*>*>{K*XN;'&[9+44[AL00>/~>?  %/ >2 )9(<(<TZZZ^=_arDI )= )KN$%0.A ((&&( x^c^g^g^i',	. #~t44  s   !Y? ?
Z	ZZ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 rS rSrg)LazyDicti  c                    Xl         / U l        g rc   dataextras)r?   rG  s     r"   rC   LazyDict.__init__  s    	r#   c                    U R                    H  u  pU R                  U" U6 -  U l        M     U R                   R                  5         U R                  $ rc   )rH  rG  clearr?   funcargss      r"   r  LazyDict.get  s@    ++JD		D$K/DI &yyr#   c                <    U R                   R                  X45        g rc   )rH  r   rL  s      r"   addLazyDict.add  s    D<(r#   rF  N)r.   rl   rm   rn   rC   r  rQ  ro   rp   r#   r"   rD  rD    s    )r#   rD  c                      \ rS rSrS rSrg)AsmDicti  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?   rV   values      r"   __missing__AsmDict.__missing__  s4    &=T']+E.455S	r#   rp   N)r.   rl   rm   rn   rY  ro   rp   r#   r"   rT  rT    s    r#   rT  c                  @   ^  \ rS rSrS rS rU 4S jrS rS rSr	U =r
$ )r  i  c           	        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>  s(     `2H$!AJJW^L_gd1gg2Hs   44cluster_dimsr  r   arch	warp_sizeKernelMetadatar  r   rp   )collectionsr]  nextr3   r	  loadsrx   r8   r
   rM   r  r=   r  r  r  pack_metadatapacked_metadatary   rW   r/   r   r   
binary_extrT  rw   r   asmkernelr   function)r?   ry   r  rW   r]  r4  r  r  re  r   r_  r`  	asm_filesrk  files                  r"   rC   CompiledKernel.__init__  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                   U R                   b  g [        R                  R                  5       n[        R                  R	                  U R
                  U R                  5      U l        [        U5      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        U l        [        R                  R+                  5       R,                  nU R                  R.                  U-  U R(                  :  a.  [        U R                  R.                  U-  U R(                  S5      eg )Nzshared memory	tmem_sizei   ztensor memorythreads)r   r   r   get_current_devicelauncher_clsry   r  runr   sharedr   rH   rs  r   load_binaryr/   rm  rn  n_regsn_spillsn_max_threadsr  rd  r   )r?   r   
max_sharedmax_tmem_sizerd  s        r"   _init_handlesCompiledKernel._init_handles  su   ;;"113==--dhhF#F+
==*, !5!5z?SS4==+..4==3J3J3VM}}&&6$T]]%<%<m_]]U[UbUbUhUhUtUtIIt{{DMM$8$8&VBRT]DK@RMM446@@	==""Y.1C1CC !8!89!DdFXFXZcdd Dr#   c                N   > US:X  a  U R                  5         [        TU ]	  U5      $ )Nrw  )r  super__getattribute__)r?   r/   	__class__s     r"   r  CompiledKernel.__getattribute__  s&    5= w'--r#   c                   [         R                  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/   rn  streamr   r   )r   runtimelaunch_enter_hookrD  r/   rn  r4   ry   r%   r*   launch_metadatar:   r6   rQ  r  )	r?   gridr  rN  retarg_dictarg_idxiarg_names	            r"   r  CompiledKernel.launch_metadata  s    ==**2		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                ~  > 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                  [        R                  R                  /	UQ76   g )Nr   r   r   )r   r   ru  get_current_streamr  rw  rn  rj  r   r  r  launch_exit_hook)r  rN  r   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]]44emm6T6T]W[]r#   )r  )r?   r  r  s   `` r"   __getitem__CompiledKernel.__getitem__  s%    !% 	] 	] r#   )rl  rn  rW   rm  r  r   r|  rz  r{  r/   rj  rw  ry   )r.   rl   rm   rn   rC   r  r  r  r  ro   __classcell__)r  s   @r"   r  r    s"    :e,.
 r#   r  )r   BaseExceptionrg   )r  r
   ri   r	   )3
__future__r   rO   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_patternr{   ptx_arg_type_patternr~   r   r%   rr   	lru_cacher   r   r   r   r   r<  r  rD  r1   rT  r  rp   r#   r"   <module>r     s   "   >  ( 6 ! . U U # #  	  	   Y 	   , 	 

# #L' 'T  1  1F O O,#$L$
 $
NN5b) ) 
d 
S Sr#   