o
    ۗiID                     @  sv  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	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ZdZeeedZ dZ!dZ"e!e!e"dZ#dd Z$d0ddZ%G dd dZ&G dd dZ'e( dd Z)d d! Z*d1d$d%Z+d2d&d'Z,d(d) Z-G d*d+ d+Z.G d,d- d-e/Z0G d.d/ d/Z1dS )3    )annotationsN   )get_cache_invalidating_env_varsir)backends)	GPUTargetAttrsDescriptor)__version__)OutOfResources)get_cache_managerget_dump_managerget_override_manager)driver)get_sass   )ast_to_ttir)Pathz^\s*tt\.func\s+(?:public\s+)?(@\w+)(\((?:%\w+: [\S\s]+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\))\s*(attributes \{[\S\s]+\})?\s+\{\s*$z=\.(?:visible|extern)\s+\.(?:entry|func)\s+(\w+)\s*\(([^)]*)\))ttirttgirptxz-%\w+: ((?:[^,\s<)]+|<[^>]+>)+(?: {[^}]+})?),?z\.param\s+\.(\w+)c                 C  sP   t d| }t d| }|d urdS t dd| } |d ur&dt|d S | S )Nz!tt\.ptr<([^,]+)ztt.nv_tma_desc = 1	nvTmaDescz {[^}]+} *r   )researchsubconvert_type_reprgroup)xmatchtma r!   V/home/ubuntu/transcripts/venv/lib/python3.10/site-packages/triton/compiler/compiler.pyr   ,   s   r   srcstrc                 C  s4   d}t || }t|dksJ dt|d }|S )Nz&"triton_gpu.num-warps"\s?=\s?(\d+)\s?:r   z(Expected exactly one match for num_warpsr   )r   findalllenint)r#   ttgir_num_warps_patternnum_warps_matches	num_warpsr!   r!   r"   _get_num_warps_from_ir_str9   s
   r+   c                   @  s0   e Zd ZddddZdd Zdd	 Zd
d ZdS )	ASTSourceNreturnNonec                 C  s   || _ d| _|j| _|| _|| _|| _t| jtr(dd t	| j
dD | _n| j D ]}t|ts8tdq-| jd u rBi | _n| j D ]}t|tsRtdqG| jd u r^t | _d S d S )Nr   c                 S  s   i | ]	\}}||  qS r!   )strip.0kvr!   r!   r"   
<dictcomp>M       z&ASTSource.__init__.<locals>.<dictcomp>,zSignature keys must be stringzConstants keys must be string)fnext__name__name	signature	constantsattrs
isinstancer$   	enumeratesplitkeys	TypeErrorr   )selfr7   r;   r<   r=   r2   r!   r!   r"   __init__E   s*   



zASTSource.__init__c                 C  sh   dd t | j D }t dd | j D }| jj d| j  d| d| }t	|
d S )Nc                 S  s   g | ]\}}|qS r!   r!   r0   r!   r!   r"   
<listcomp>\   s    z"ASTSource.hash.<locals>.<listcomp>c                 s  s     | ]\}}t ||fV  qd S N)r$   r0   r!   r!   r"   	<genexpr>_   s    z!ASTSource.hash.<locals>.<genexpr>-utf-8)sortedr;   itemsr<   r7   	cache_keyr=   hashhashlibsha256encode	hexdigest)rC   
sorted_sigsorted_constantskeyr!   r!   r"   rM   [   s   $zASTSource.hashc                 C  s   t | j| ||||dS )N)contextoptionscodegen_fns
module_map)r   r7   )rC   rV   rW   rX   rU   r!   r!   r"   make_irc   s   zASTSource.make_irc                 C  s   t  S rF   )dictrC   r!   r!   r"   parse_optionsg   s   zASTSource.parse_optionsNNr-   r.   r9   
__module____qualname__rD   rM   rY   r\   r!   r!   r!   r"   r,   C   s
    r,   c                   @  s,   e Zd Zdd Zdd Zdd Zdd Zd	S )
IRSourcec                 C  s   || _ t|}|jdd  | _| | _tt| j | jtj	}|
d| _|
d}tt| j |}dd t|D | _d S )Nr   r   c                 S  s   i | ]	\}}|t |qS r!   )r   )r1   r2   tyr!   r!   r"   r4   v   r5   z%IRSource.__init__.<locals>.<dictcomp>)pathr   suffixr8   	read_textr#   r   r   prototype_pattern	MULTILINEr   r:   r%   arg_type_patternr?   r;   )rC   rd   r   r;   typesr!   r!   r"   rD   m   s   

zIRSource.__init__c                 C  s   t | jd S )NrI   )rN   rO   r#   rP   rQ   r[   r!   r!   r"   rM   x   s   zIRSource.hashc                 C  s   t | j|}||_|S rF   )r   parse_mlir_modulerd   rU   )rC   rV   rW   rX   rU   moduler!   r!   r"   rY   {   s   zIRSource.make_irc                 C  s   | j dkrdt| jiS t S )Nr   r*   )r8   r+   r#   rZ   r[   r!   r!   r"   r\      s   
zIRSource.parse_optionsNr_   r!   r!   r!   r"   rb   k   s
    rb   c               
   C  s  dd l } tjtjtjt}g }ttd}|t|	 
 g7 }W d    n1 s0w   Y  tj|ddftj|ddfg}|D ]6\}}| j|g|dD ])}t|j|jjd}|t|	 
 g7 }W d    n1 syw   Y  qUqIt }ttj|dd}	 |	d
}	|	sn||	 qW d    n1 sw   Y  ||
  tj|d}
| j|
gddD ])}t|j|jjd}|t|	 
 g7 }W d    n1 sw   Y  qt d| S )Nr   rbcompilerztriton.compiler.r   ztriton.backends.)prefixz_C/libtriton.soTi   languageztriton.language.rH   )pkgutilosrd   dirnameabspath__file__openrN   rO   readrQ   joinwalk_packagesmodule_finder	find_specr:   originupdateappendr	   )rq   TRITON_PATHcontentsfpath_prefixesrd   ro   liblibtriton_hashchunklanguage_pathr!   r!   r"   
triton_key   sD   

r   c                 C  sZ   |dks|dkrt | |}||_|S |dks|dkr!t|  S |dkr+t|  S d S )Nr   r   llirr   cubin)r   rk   rU   r   rf   
read_bytes)	full_namer8   rU   rl   r!   r!   r"   parse   s   r   eBaseExceptionc                   s   t dddkr
dS | jdurt| j | jdurt| j ddg}| j g } durBt fdd|D s;|   j  dus+t	||d	d D ]\}}||_qK|sZd| _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!"
    TRITON_FRONT_END_DEBUGGING01Nz"/triton/compiler/code_generator.pyz/ast.pyc                 3  s$    | ]} j jj|r|V  qd S rF   )tb_framef_codeco_filenameendswith)r1   r   tbr!   r"   rG      s   " z#filter_traceback.<locals>.<genexpr>r   r   )
rr   getenv	__cause__filter_traceback__context____traceback__anyr~   tb_nextzip)r   	BAD_FILESframes	cur_frame
next_framer!   r   r"   r      s,   






r   c           "      C  sL  |d u r	t j }t|tsJ dt|}t| t }|r+t| ts'J dt| } | 	 }|	t
|p6t
 fi |}t }t  d|   d|  d|  dtt|  	}t|d }t|}	tjdddk}
tjdddk}|
rt|  nd }|rt|  nd }| jd d	 }| d
}|	|pi }||}tjdddk}|s|d urtt| }t | ||S ||d|j!|}t
 }|"|| t#|$ %| j&}|r|d7 }t'( }t')| |)| |* }|+ }z
| ,||||}W n t-y } zt.|  d }~ww tjdd }t#| |d  D ]X\}}|||}| d| }|d urZ|/| } d urZt0d|   t1| ||}|	2||||< |d urm|2|| ||kr|	/|}!|3|! t0d|!  |}q.|	j2tj4|t5d|dd||< |	6|| |7  t | ||S )Nz target must be of GPUTarget typez'source must be either AST or a filepathrH   rI   TRITON_KERNEL_OVERRIDEr   r   TRITON_KERNEL_DUMP   .jsonTRITON_ALWAYS_COMPILE)rM   targetr   
USE_IR_LOC.z
Overriding kernel with file zCreating new locations for )defaultF)binary)8r   activeget_current_targetr>   r   make_backendr,   r$   rb   r\   rZ   r   r   rM   rJ   rK   rN   rO   rP   rQ   r   rr   environgetr   r   r:   	get_groupjsonloadsr   rf   CompiledKernel__dict__
add_stageslistrA   indexr8   r   rU   load_dialectsget_codegen_implementationget_module_maprY   	Exceptionr   get_fileprintr   putcreate_location_snapshotdumpsvars	put_groupdisable_multithreading)"r#   r   rV   backend	ir_sourceextra_optionsenv_varsrT   rM   fn_cache_managerenable_overrideenable_ir_dumpfn_override_managerfn_dump_manager	file_namemetadata_filenamemetadata_groupmetadata_pathalways_compilemetadatastagesfirst_stagerU   rW   rX   rl   r   
use_ir_locr8   
compile_irnext_moduleir_filenamer   ir_full_namer!   r!   r"   compile   s   
:









r   c                   sN    fddt  D }t|dkr!tt| d j d| d|d  S )Nc                   s   g | ]}|j  r|j qS r!   )rn   supports_target)r1   r   r   r!   r"   rE   3  s    z make_backend.<locals>.<listcomp>r   z! compatible backends for target (z) (z). There should only be one.r   )r   valuesr&   RuntimeErrorr   )r   activesr!   r   r"   r   2  s   r   c                   @  s&   e Zd Zdd Zd
ddZdd Zd	S )LazyDictc                 C  s   || _ g | _d S rF   )dataextras)rC   r   r!   r!   r"   rD   <  s   
zLazyDict.__init__r-   r.   c                 C  s0   | j D ]\}}| j|| B | _q| j   | jS rF   )r   r   clearrC   funcargsr!   r!   r"   r   @  s   
zLazyDict.getc                 C  s   | j ||f d S rF   )r   r~   r   r!   r!   r"   addF  s   zLazyDict.addNr^   )r9   r`   ra   rD   r   r   r!   r!   r!   r"   r   :  s    
r   c                   @  s   e Zd Zdd ZdS )AsmDictc                 C  s.   |dkrt | d }ntd| || |< |S )Nsassr   zUnknown key: '%s')r   KeyError)rC   rT   valuer!   r!   r"   __missing__L  s
   zAsmDict.__missing__N)r9   r`   ra   r   r!   r!   r!   r"   r   J  s    r   c                      sD   e Zd ZdZdZdd Zdd Z fddZdd	 Zd
d Z	  Z
S )r   Nc                   s  ddl m} tdd | D }t| }t|d |d< |d }t|d |d |d	 |d< |d
t	t
| }|di || _t| jj}	|	| j| _|| _|| _| jj| _dd | D }
|	j t fdd|
D | _| j  | _d | _d | _d S )Nr   )
namedtuplec                 s  s&    | ]\}}| d rt|V  qdS )r   Nr   r   r1   cpr!   r!   r"   rG   `  s   $ z*CompiledKernel.__init__.<locals>.<genexpr>cluster_dimsr   r   arch	warp_sizeKernelMetadatac                 S  s"   g | ]\}}| d st|qS )r   r   r   r!   r!   r"   rE   n  s   " z+CompiledKernel.__init__.<locals>.<listcomp>c                   s:   i | ]}|j d d |j d d  kr| n| qS )r   N)re   r   rf   )r1   file
binary_extr!   r"   r4   p  s    ,z+CompiledKernel.__init__.<locals>.<dictcomp>r!   )collectionsr   nextrK   r   r   rf   tupler   rJ   r   rA   r   r   r   pack_metadatapacked_metadatar#   rM   r:   r   r   asmkernelrl   function)rC   r#   r   rM   r   r   r   r   r   r   	asm_filesr!   r   r"   rD   ^  s*   


zCompiledKernel.__init__c                 C  s   | j d urd S tj }tj| j| j| _tjj	|d }| jj
|kr-t| jj
|dtjj| j| j| jj
|\| _ | _| _| _d S )Nmax_shared_memzshared memory)rl   r   r   get_current_devicelauncher_clsr#   r   runutilsget_device_propertiessharedr
   load_binaryr:   r  r  n_regsn_spills)rC   device
max_sharedr!   r!   r"   _init_handles{  s   

zCompiledKernel._init_handlesc                   s   |dkr|    t |S )Nr	  )r  super__getattribute__)rC   r:   	__class__r!   r"   r    s   zCompiledKernel.__getattribute__c           	      G  s   t jd u rd S t| j| j|d}t| jtr| jjj	d u r |S i }d}t
| jjjD ]\}}|| jjjv r?| jj| ||< q+|| ||< |d7 }q+|| jjj	|| j|f |S )N)r:   r  streamr   r   )r   launch_enter_hookr   r:   r  r>   r#   r,   r7   launch_metadatar?   	arg_names
constexprsr<   r   r   )	rC   gridr  r   retarg_dictarg_idxiarg_namer!   r!   r"   r    s   

zCompiledKernel.launch_metadatac                   s       d d fdd
}|S )N)r  c              
     sl   | d u rt j }t j|} j | g|R  }j d  d  d | jj|tj	tj
g	|R   d S )Nr   r   r   )r   r   r  get_current_streamr  r	  r  r  r   r  launch_exit_hook)r  r   r  r  r  rC   r!   r"   runner  s   
"z*CompiledKernel.__getitem__.<locals>.runner)r  )rC   r  r%  r!   r$  r"   __getitem__  s   zCompiledKernel.__getitem__)r9   r`   ra   r  r#  rD   r  r  r  r&  __classcell__r!   r!   r  r"   r   W  s    r   )r#   r$   )r   r   r]   )2
__future__r   rN   r   _C.libtritonr   r   r   backends.compilerr   r   r   r	   runtime.autotunerr
   runtime.cacher   r   r   runtime.driverr   tools.disasmr   code_generatorr   pathlibr   r   	functoolsrr   mlir_prototype_patternptx_prototype_patternrg   mlir_arg_type_patternptx_arg_type_patternri   r   r+   r,   rb   	lru_cacher   r   r   r   r   r   rZ   r   r   r!   r!   r!   r"   <module>   sR    


(
"

%Y