
    Xh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j'                    d             Z( ej'                    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)          Z2dS )-    )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|           }|dS t          j        dd|           } |%dt          |                    d                    z   S | S )Nz!tt\.ptr<([^,]+)ztt.nv_tma_desc = 1	nvTmaDescz {[^}]+} *   )researchsubconvert_type_reprgroup)xmatchtmas      j/var/www/tools.fuzzalab.pt/emblema-extractor/venv/lib/python3.11/site-packages/triton/compiler/compiler.pyr   r   '   sn     I)1--E
))1
-
-C
{
{B""A&u{{1~~6666H    c                  *    e Zd Zdd	dZd Zd Zd ZdS )
	ASTSourceNreturnNonec                   || _         t          j        | _        d| _        |j        | _        || _        t                      | _	        |m|
                                D ]X\  }}t          |t                    r|j                            |          fn|}t          |t                    sJ || j	        |<   Y|pt                      | _        t          | j        t                    r8d t#          | j                            d                    D             | _        d S | j                                        D ]&}t          |t                    st)          d          'd S )Nttirc                >    i | ]\  }}||                                 S  )strip.0kvs      r"   
<dictcomp>z&ASTSource.__init__.<locals>.<dictcomp>D   s&    \\\tq!a\\\r#   ,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keys	TypeError)selfr3   r9   
constexprsrB   r/   r0   s          r"   __init__zASTSource.__init__6   sS    K	"!"((** & &11;As1C1CJR\''**--!!U+++++$%q!!_dff
dnc** 	E\\yAUAUVYAZAZ7[7[\\\DNNN^((** E E!!S)) E#$CDDDEE Er#   c                   d t          | j                                                  D             }d d                    fdt          | j                                                  D                       }| j        j         dt          | j                   d| d| }t          j
        |                    d                                                    S )Nc                    g | ]\  }}|S r+   r+   r-   s      r"   
<listcomp>z"ASTSource.hash.<locals>.<listcomp>K   s    CCCDAqaCCCr#   c                N    t          | d          r| j        nt          |           S )N	cache_key)hasattrrN   r>   )r   s    r"   <lambda>z ASTSource.hash.<locals>.<lambda>L   s!    71k+B+BNAKKA r#   -c                ,    g | ]\  }} |          S r+   r+   )r.   r/   r0   get_keys      r"   rL   z"ASTSource.hash.<locals>.<listcomp>M   s%    !X!X!XA''!**!X!X!Xr#   utf-8)sortedr9   r<   joinr;   r3   rN   r>   rB   hashlibsha256encode	hexdigest)rG   
sorted_sigconstants_keykeyrS   s       @r"   hashzASTSource.hashJ   s    CCF4>+?+?+A+A$B$BCCC
NN!X!X!X!X@T@T@V@V9W9W!X!X!XYY"SSS__SSzSSMSS~cjj1122<<>>>r#   c                :    ddl m}  || j        | ||||          S )Nr   )ast_to_ttir)contextoptionscodegen_fns
module_map)code_generatorr`   r3   )rG   rb   rc   rd   ra   r`   s         r"   make_irzASTSource.make_irQ   s=    //////{47D'7Xc&02 2 2 	2r#   c                    t                      S N)r:   rG   s    r"   parse_optionszASTSource.parse_optionsV   s    vvr#   NNr&   r'   r7   
__module____qualname__rI   r^   rf   rj   r+   r#   r"   r%   r%   4   s_        E E E E E(? ? ?2 2 2
    r#   r%   c                  &    e Zd Zd Zd Zd Zd ZdS )IRSourcec                   || _         t          |          }|j        dd          | _        t          j        | _        |                                | _        t          j
        |           |
                    |           | j        dk    rt          j        t          | j                 | j        t          j                  }|                    d          | _        |                    d          }t          j        t$          | j                 |          }d t'          |          D             | _        d S t          j        | j         |          | _        | j                                        }d|z   | _        | j                            |          }| j                            |          }	d t'          |	          D             | _        d S )Nr   r   r   c                4    i | ]\  }}|t          |          S r+   )r   r.   r/   tys      r"   r1   z%IRSource.__init__.<locals>.<dictcomp>l   s'    UUU51ba!22!6!6UUUr#   @c                    i | ]\  }}||	S r+   r+   rt   s      r"   r1   z%IRSource.__init__.<locals>.<dictcomp>s   s    DDD2aDDDr#   )pathr   suffixr6   r   r4   r5   	read_textsrcr   load_dialectsr   r   prototype_pattern	MULTILINEr   r8   findallarg_type_patternrC   r9   parse_mlir_modulemoduleget_entry_func_nameget_functionget_function_signature)
rG   rx   ra   backendr    r9   typesfn_namefuncOpfunc_tys
             r"   rI   zIRSource.__init__\   sk   	Dzz;qrr? >>##
!!!g&&& 8uI/948R\RREADIAIJ/99EEEUUIeDTDTUUUDNNN.ty'BBDKk5577GgDI[--g66Fk88@@GDD71C1CDDDDNNNr#   c                ~    t          j        | j                            d                                                    S )NrT   )rW   rX   r{   rY   rZ   ri   s    r"   r^   zIRSource.hashu   s,    ~dhoog6677AACCCr#   c                (    || j         _        | j         S rh   )r   ra   )rG   rb   rc   rd   ra   s        r"   rf   zIRSource.make_irx   s    %{r#   c                    | j         dk    r*| j                            d          }|
J d            d|iS t                      S )Nttgirzttg.num-warpsz'Unable to parse ttg.num-warps attribute	num_warps)r6   r   get_int_attrr:   )rG   r   s     r"   rj   zIRSource.parse_options|   sM    8w00AAI((*S(((++vvr#   Nrm   r+   r#   r"   rq   rq   Z   sV        E E E2D D D      r#   rq   c                    dd l } t          j                            t          j                            t          j                            t
                                        }g }t          t
          d          5 }|t          j        |	                                          
                                gz  }d d d            n# 1 swxY w Y   t          j                            |d          dft          j                            |d          dfg}|D ]\  }}|                     |g|          D ]}t          |j                            |j                  j        d          5 }|t          j        |	                                          
                                gz  }d d d            n# 1 swxY w Y   t          j                    }t#          j        d                              d	          d
         }	t          t          j                            |dd|	           d          5 }	 |	                    d          }
|
sn|                    |
           .	 d d d            n# 1 swxY w Y   |                    |
                                           t          j                            |d          }|                     |gd          D ]}t          |j                            |j                  j        d          5 }|t          j        |	                                          
                                gz  }d d d            n# 1 swxY w Y   t,           d                    |          z   S )Nr   rbcompilerztriton.compiler.r   ztriton.backends.)prefix
EXT_SUFFIX._Cz
libtriton.Ti   r5   ztriton.language.rQ   )pkgutilosrx   dirnameabspath__file__openrW   rX   readrZ   rV   walk_packagesmodule_finder	find_specr8   origin	sysconfigget_config_varrD   updateappendr   )r   TRITON_PATHcontentsfpath_prefixesrx   r   liblibtriton_hashr6   chunklanguage_paths               r"   
triton_keyr      s   NNN'//"'//"'//(2K2K"L"LMMKH	h		 ;W^AFFHH--7799::; ; ; ; ; ; ; ; ; ; ; ; ; ; ; 
k:	.	.0BC	k:	.	.0BCM & C Cf(($(?? 	C 	CCc'11#(;;BDII CQW^AFFHH55??AABBC C C C C C C C C C C C C C C	C
 ^%%N

"<
0
0
6
6s
;
;B
?C	bgll;.@3.@.@AA4	H	H )A	)FF7OOE !!%(((		) 	) ) ) ) ) ) ) ) ) ) ) ) ) ) ) OON,,..///GLLj99M$$m_=O$PP ? ?##--ch77>EE 	?11;;==>>H	? 	? 	? 	? 	? 	? 	? 	? 	? 	? 	? 	? 	? 	? 	?chhx0000sH   :=CC
C%=F..F25F2-0I**I.1I.=MM	M	c                V    t           j        j                            |           d         S )Nmax_shared_mem)r   activeutilsget_device_properties)devices    r"   r   r      s!    =44V<<=MNNr#   c                   |dk    s|dk    rt          j        | |          }||_        |S |dk    s|dk    s|dk    r!t          |                                           S |dk    s|dk    r!t          |                                           S d S )Nr)   r   llirr   amdgcncubinhsaco)r   r   ra   r   rz   
read_bytes)	full_namer6   ra   r   s       r"   parser      s    
f}}w%i99 
f}}uxI((***
g~~I))+++ (r#   eBaseExceptionc                   t           j        j        rdS | j        t	          | j                   | j        t	          | j                   ddg}d |D             }| j        g }9t          fd|D                       s|                               j	        9t          ||dd                   D ]\  }}||_	        |s	d| _        dS d|d         _	        |d         | _        dS )	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.pyc                N    g | ]"}|                     d t          j                  #S )/)replacer   sep)r.   bad_files     r"   rL   z$filter_traceback.<locals>.<listcomp>   s*    III8!!#rv..IIIr#   c              3  d   K   | ]*}j         j        j                            |          &|V  +d S rh   )tb_framef_codeco_filenameendswith)r.   r   tbs     r"   	<genexpr>z#filter_traceback.<locals>.<genexpr>   s?      VV2;+=+I+R+RST+U+UV1VVVVVVr#   r   r   r   )r   compilationfront_end_debugging	__cause__filter_traceback__context____traceback__anyr   tb_nextzip)r   	BAD_FILESframes	cur_frame
next_framer   s        @r"   r   r      s$    , {%%%} ''' 	-I JIyIIII	
BF
.VVVViVVVVV 	MM"Z .
 $'vvabbz#:#: ' 'J&	 $!r
 )r#   c                  .    e Zd ZddZddZddZdd	Zd
S )CompileTimerr&   r'   c                `    t          j                     | _        d | _        g | _        d | _        d S rh   )timestartir_initialization_endlowering_stage_endsstore_results_endri   s    r"   rI   zCompileTimer.__init__   s-     IKK
37"<> /3r#   c                6    t          j                     | _        d S rh   )r   r   ri   s    r"   finished_ir_initializationz'CompileTimer.finished_ir_initialization   s    %)Y[["""r#   
stage_namer>   c                `    | j                             |t          j                    f           d S rh   )r   r   r   )rG   r   s     r"   stage_finishedzCompileTimer.stage_finished   s*     ''TY[[(ABBBBBr#   knobs.CompileTimesc                B   t          j                     }| j        || _        n|| _        d	d}g }| j        }| j        D ](\  }}|                    | |||          f           |})t          j         || j        | j                  | ||| j                            S )
Nr   floatendfloat | Noner&   intc                4    |dS t          || z
  dz            S )Nr   i@B )r   )r   r   s     r"   deltazCompileTimer.end.<locals>.delta   s#    {qew.///r#   )ir_initializationlowering_stagesstore_results)r   r   r   r   r&   r   )r   r   r   r   r   r   CompileTimesr   )rG   	timestampr   lowering_stage_durationsstage_startr   	stage_ends          r"   r   zCompileTimer.end   s    IKK	%-)2D&&%.D"	0 	0 	0 	0
 $& 0%)%= 	$ 	$!J	$++Z{I9V9V,WXXX#KK!#eDJ0JKK4%T-CDD
 
 
 	
r#   Nrl   )r   r>   r&   r'   )r&   r   )r7   rn   ro   rI   r   r   r   r+   r#   r"   r   r      sh        4 4 4 41 1 1 1C C C C
 
 
 
 
 
r#   r   c                   t           j        j        }|rt                      }|t          j                                        }t          |t                    s
J d            t          |          }t          | t                     }|rCt          | t                    s
J d            t          j                    }t          | ||          } |                                 }|                    t!          |pt!                      fi |          }t#                      }	t%                       d|                                  d|                                 d|                                 dt          t)          |	                                                     	}
t-          j        |
                    d                                                    }t5          |          }t           j        j        }t           j        j        }t           j        j        }|r!t=          |                                           nd }|r!t?          |                                           nd }| j         d d         }| d}|!                    |          pi }|"                    |          }t           j        j#        }|sP|NtI          | ||          }|r9 || |j%        &                                ||'                                d           |S ||d	|j(        |	}tR          |d
<   t!                      }|*                    ||| j+                   tY          |-                                          .                    | j/                  }|r|dz  }t          | t                    s<t          j                    }t          j0        |           |0                    |           |1                    |          }|2                                }	 | 3                    ||||          }n"# th          $ r}tk          |            d }~ww xY w|r&| d| j/         }|6                    ||          ||<   n| d}|6                    ||          ||<   t           j        j7        } |r3| r1|8                    | j9                   tu          d| j9                    |r|;                                 tY          |                                          |d          D ])\  }!}" |"||          }#| d|! }|B|"                    dd           x}$r)|$<                    d|!           rt{          |$|!|          }#n:|>                    |          x}%r#tu          d|%            t{          |%|!|          }#|r|!dv r|6                    |#|          ||<   ||6                    |#|           | |!k    r<|>                    |          }&|#8                    |&           tu          d|&            |#}|r|?                    |!           +|6                    t          jA        |t                    |d          ||<   |C                    ||           t           j        jD        s|E                                 |r" || |||'                                d           tI          | ||          S )Nz target must be of GPUTarget typez'source must be either AST or a filepathrQ   rT      .jsonT)r{   metadatametadata_grouptimes	cache_hit)r^   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_targetr=   r
   make_backendr%   r>   r   ra   rq   rj   r:   r   r   r^   rU   r<   rW   rX   rY   rZ   r   overridedump_irstore_binary_onlyr   r   r8   	get_groupgetalways_compileCompiledKernelr   _asdictr   __dict__r   
add_stagesr5   listrE   r@   r6   r|   get_codegen_implementationget_module_maprf   	Exceptionr   put
use_ir_loccreate_location_snapshotrx   printr   r   r   get_filer   r  dumpsvars	put_groupenable_asandisable_multithreading)'r{   r   rb   compilation_listenertimerr   	ir_sourcera   extra_optionsenv_varsr]   r^   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_stagerc   rd   r   r   ir_filenamer  r6   
compile_irnext_moduler  r   ir_full_names'                                          r"   compiler5    s2    ,5 ~1133fi((LL*LLLL6""GsI...I .#s##NN%NNNN*,,sGW--%%''M##D):DFF$L$Lm$L$LMMG.00H\\
j
jCHHJJ
j
j
j
j',,..
j
j3vV^VdVdVfVfOgOgKhKh
j
jC>#**W--..88::D(.. '0O&.N);>MW.sxxzz:::SW6DN&sxxzz222$O
 #I$+++%//0ABBHbN"&&'899M&5N m7S.$77 	  --//-iikk    
   
 	H "-HVVFvw555v{{}}%%++CG44K q c8$$ '*,,
!!!g&&&44W==K''))JWk:wGG     P"..SW..&6&:&:6;&O&O{##"+++&6&:&:6;&O&O{#"-J 8Z 8''1116CH66777 +((***//= & &Z j22"**S**&  (||M4@@@ ?kFZFZ[d_b[d[dFeFe ?#Kg>>-66{CCCY 	9>9>>???	388K! 	Ys.H'H'H*:*>*>{K*X*XN;'&[999+44[AAL00>>>>>>??? 	&  %%%(8(<(<TZZ^=_=_=_arDI )= )K )KN$%0.AAA ( )&&(((  .x^c^g^g^i^i',	. 	. 	. 	. #~t444s   >P 
P6!P11P6r   r
   r&   r	   c                      fdt          j                    D             }t          |          dk    r*t          t          |           d j         d| d           |d                    S )Nc                R    g | ]#}|j                                       |j         $S r+   )r   supports_target)r.   r   r   s     r"   rL   z make_backend.<locals>.<listcomp>  s1    [[[a
8R8RSY8Z8Z[qz[[[r#   r   z! compatible backends for target (z) (z). There should only be one.r   )r   valueslenRuntimeErrorr   )r   activess   ` r"   r  r    s    [[[[8?#4#4[[[G
7||q7||vvfnvvQXvvvx x 	x71:fr#   c                       e Zd Zd Zd Zd ZdS )LazyDictc                "    || _         g | _        d S rh   )dataextras)rG   r@  s     r"   rI   zLazyDict.__init__  s    	r#   c                    | j         D ]\  }}| j         || z  | _        | j                                          | j        S rh   )rA  r@  clearrG   funcargss      r"   r  zLazyDict.get  sI    + 	0 	0JD$	DD$K/DIIyr#   c                >    | j                             ||f           d S rh   )rA  r   rD  s      r"   addzLazyDict.add  s"    D$<(((((r#   N)r7   rn   ro   rI   r  rH  r+   r#   r"   r>  r>    sA            ) ) ) ) )r#   r>  c                      e Zd Zd ZdS )AsmDictc                l    |dk    rt          | d                   }nt          d|z            || |<   |S )Nsassr   zUnknown key: '%s')r   KeyError)rG   r]   values      r"   __missing__zAsmDict.__missing__  s@    &==T']++EE.4555S	r#   N)r7   rn   ro   rO  r+   r#   r"   rJ  rJ    s#            r#   rJ  c                  6     e Zd Zd Zd Z fdZd Zd Z xZS )r  c           	     D   ddl m} t          d |                                D                       }t	          j        |                                          }t          |d                   |d<   |d         }t          |d         |d         |d                   |d<    |d	t          t          |                                                              } |di || _        t          | j        j                  }	|	                    | j                  | _        || _        || _        | j        j        | _        d
 |                                D             }
|	j        t+          fd|
D                       | _        | j                 | _        d | _        d | _        d S )Nr   )
namedtuplec              3  d   K   | ]+\  }}|                     d           t          |          V  ,dS )r   Nr   r   r.   cps      r"   r   z*CompiledKernel.__init__.<locals>.<genexpr>  s>      ``$!QAJJW^L_L_`d1gg``````r#   cluster_dimsr   r   arch	warp_sizeKernelMetadatac                \    g | ])\  }}|                     d           t          |          *S )r   rT  rU  s      r"   rL   z+CompiledKernel.__init__.<locals>.<listcomp>  s4    [[[AqzzRYGZGZ[T!WW[[[r#   c                    i | ]K}|j         d d         |j         d d         k    r|                                n|                                LS )r   N)ry   r   rz   )r.   file
binary_exts     r"   r1   z+CompiledKernel.__init__.<locals>.<dictcomp>  se     
 
 
 KO$+abb/Z2O2OT__...UYUcUcUeUe
 
 
r#   r+   )collectionsrR  nextr<   r  loadsrz   rA   r
   rU   r  rE   r   r  r   pack_metadatapacked_metadatar{   r^   r8   r_  rJ  asmkernelr   function)rG   r{   r   r^   rR  r-  r   r   r[  r   	asm_filesr_  s              @r"   rI   zCompiledKernel.__init__  s   ******``.2F2F2H2H```aa:m557788#(.)A#B#B (#&vi'8&.&Q\J]^^#$4fT(--//=R=R6S6STT&2222t}344&44T]CC	M&	[[)=)=)?)?[[[	'
 
 
 
 
!
 
 
   hz* r#   c                2   | j         d S t          j                                        }t          j                            | j        | j                  | _        t          |          }| j        j	        |k    rt          | j        j	        |d          t          | j        d          r9| j        j        -d}| j        j        |k    rt          | j        j        |d          t          j        j                            | j        | j        | j        j	        |          \  | _         | _        | _        | _        | _        t          j                                        j        }| j        j        |z  | j        k    r#t          | j        j        |z  | j        d          d S )Nzshared memory	tmem_sizei   ztensor memorythreads)r   r   r   get_current_devicelauncher_clsr{   r   runr   sharedr   rO   rj  r   load_binaryr8   rf  rg  n_regsn_spillsn_max_threadsr  rZ  r   )rG   r   
max_sharedmax_tmem_sizerZ  s        r"   _init_handleszCompiledKernel._init_handles  s\   ;"F1133=--dhFF#F++
=*,, !5z?SSS4=+.. 	^4=3J3VM}&66$T]%<m_]]]U[UbUhUtUtIt{DM$8&VB VBRT]DK@RM4466@	="Y.1CCC !89!DdFXZcddd DCr#   c                z    |dk    r|                                   t                                          |          S )Nrn  )rv  super__getattribute__)rG   r8   	__class__s     r"   ry  zCompiledKernel.__getattribute__  s6    5==   ww''---r#   c                   t           j        j        d S 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        |f           |S )N)r8   rg  streamr   r   )r   runtimelaunch_enter_hookr>  r8   rg  r=   r{   r%   r3   launch_metadatarC   r?   rH  r   )	rG   gridr|  rF  retarg_dictarg_idxiarg_names	            r"   r  zCompiledKernel.launch_metadata  s    =*24	t}PVWWXX$(I.. 	$(+2M2UJ$TX[%:;; 	 	KAx!%gHXqLGG+dDM8-LMMM
r#   c                D                                        d d fd
}|S )N)r|  c                D   | =t           j                                        }t           j                            |          }  j        | g|R  } j        d         d         d         | j        j        |t          j	        j
        t          j	        j        g	|R   d S )Nr   r   r   )r   r   rl  get_current_streamr  rn  rg  rd  r   r}  r~  launch_exit_hook)r|  rF  r   r  r  rG   s       r"   runnerz*CompiledKernel.__getitem__.<locals>.runner  s    ~99;;99&AA2d24G$GGGODHT!Wd1gtAwtG[]l]4em6T]W[] ] ] ] ] ]r#   )rv  )rG   r  r  s   `` r"   __getitem__zCompiledKernel.__getitem__  sP    !% 	] 	] 	] 	] 	] 	] 	] 	] r#   )	r7   rn   ro   rI   rv  ry  r  r  __classcell__)rz  s   @r"   r  r    sw          :e e e,. . . . .
        r#   r  )r   r   rk   )r   r
   r&   r	   )3
__future__r   rW   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%   rq   	lru_cacher   r   r   r   r   r5  r  r>  r:   rJ  r  r+   r#   r"   <module>r     s   " " " " " "   > > > > > > > >       ( ( ( ( ( ( 6 6 6 6 6 6 6 6 ! ! ! ! ! ! ! ! . . . . . . U U U U U U U U U U # # # # # # # # # # # #       				     				      Y 	   , 	 

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