o
    oh&H                     @  s`  d dl mZ d dlZd dlZddlmZmZ ddlmZ ddlm	Z	 ddl
mZ ddlmZ dd	lmZmZmZ dd
lmZ ddlmZ ddlmZ d dlmZ d dlZd dlZd dlZd dlZdZdeiZdZ de iZ!dd Z"G dd dZ#G dd dZ$e% dd Z&dd Z'd*ddZ(d+d d!Z)d"d# Z*G d$d% d%Z+G d&d' d'e,Z-G d(d) d)Z.dS ),    )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                 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   l/var/www/html/construction_image-detection-poc/venv/lib/python3.10/site-packages/triton/compiler/compiler.pyr   '   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 | _|d ur8| D ]\}}t|t	r)|j
|fn|}t|ts2J || j|< q|p<t | _t| jt	rTdd t| jdD | _d S | j D ]}t|t	sdtdqYd S )Nttirc                 S  s   i | ]	\}}||  qS r   )strip.0kvr   r   r   
<dictcomp>C       z&ASTSource.__init__.<locals>.<dictcomp>,zSignature keys must be string)fnext__name__name	signaturedict	constantsitems
isinstancestr	arg_namesindextupleattrs	enumeratesplitkeys	TypeError)selfr,   r0   
constexprsr9   r'   r(   r   r   r   __init__6   s$    
zASTSource.__init__c                   sz   dd t | j D }dd  d fddt | j D }| jj dt| j d| d| }t	
|d S )Nc                 S  s   g | ]\}}|qS r   r   r%   r   r   r   
<listcomp>J   s    z"ASTSource.hash.<locals>.<listcomp>c                 S  s   t | dr| jS t| S )N	cache_key)hasattrrB   r5   )r   r   r   r   <lambda>K       z ASTSource.hash.<locals>.<lambda>-c                   s   g | ]\}} |qS r   r   r%   get_keyr   r   rA   L   rE   utf-8)sortedr0   r3   joinr2   r,   rB   r5   r9   hashlibsha256encode	hexdigest)r>   
sorted_sigconstants_keykeyr   rG   r   hashI   s
   "$zASTSource.hashc                 C  s   t | j| ||||dS )N)contextoptionscodegen_fns
module_map)r   r,   r>   rU   rV   rW   rT   r   r   r   make_irP   s   zASTSource.make_irc                 C  s   t  S N)r1   r>   r   r   r   parse_optionsT   s   zASTSource.parse_optionsNNr!   r"   r.   
__module____qualname__r@   rS   rY   r\   r   r   r   r   r    4   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| || | jdkrOt	t
| j | jtj}|d| _|d}tt| j |}dd t|D | _d S t| j || _| j }d| | _| j|}| j|}	dd t|	D | _d S )Nr   r   r   c                 S  s   i | ]	\}}|t |qS r   )r   r&   r'   tyr   r   r   r)   i   r*   z%IRSource.__init__.<locals>.<dictcomp>@c                 S  s   i | ]\}}||qS r   r   rc   r   r   r   r)   p   s    )pathr   suffixr-   	read_textsrcr   load_dialectsr   r   prototype_pattern	MULTILINEr   r/   findallarg_type_patternr:   r0   parse_mlir_modulemoduleget_entry_func_nameget_functionget_function_signature)
r>   rf   rT   backendr   r0   typesfn_namefuncOpfunc_tyr   r   r   r@   Z   s$   






zIRSource.__init__c                 C  s   t | jd S )NrI   )rL   rM   ri   rN   rO   r[   r   r   r   rS   r   s   zIRSource.hashc                 C  s   || j _| j S rZ   )rp   rT   rX   r   r   r   rY   u   s   zIRSource.make_irc                 C  s4   | j dkr| jd}|d usJ dd|iS t S )Nttgirzttg.num-warpsz'Unable to parse ttg.num-warps attribute	num_warps)r-   rp   get_int_attrr1   )r>   rz   r   r   r   r\   y   s
   
zIRSource.parse_optionsNr_   r   r   r   r   rb   X   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dd	d
 }	ttj|d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.)prefix
EXT_SUFFIX._Cz
libtriton.Ti   languageztriton.language.rF   )pkgutilosrf   dirnameabspath__file__openrL   rM   readrO   rK   walk_packagesmodule_finder	find_specr/   origin	sysconfigget_config_varr;   updateappendr   )r   TRITON_PATHcontentsfpath_prefixesrf   r~   liblibtriton_hashr-   chunklanguage_pathr   r   r   
triton_key   sF   

r   c                 C  sj   |dks|dkrt | |}||_|S |dks|dks|dkr%t|  S |dks-|dkr3t|  S d S )Nr#   ry   llirr   amdgcncubinhsaco)r   ro   rT   r   rh   
read_bytes)	full_namer-   rT   rp   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}dd |D }| j g } durIt fd	d
|D sB|   j  dus2t	||dd D ]\}}||_qR|sad| _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                 S  s   g | ]	}| d tjqS )/)replacer   sep)r&   bad_filer   r   r   rA      r*   z$filter_traceback.<locals>.<listcomp>c                 3  s$    | ]} j jj|r|V  qd S rZ   )tb_framef_codeco_filenameendswith)r&   r   tbr   r   	<genexpr>   s   " z#filter_traceback.<locals>.<genexpr>r   r   r   )
r   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  s  |d u r	t j }t|tsJ dt|}t| t }|r1t| ts'J dt	 }t
| ||} |  }|t|p<t 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}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| ||	S |	|d|j|}t |d< t }|!|| t"|# $| j%}|r|d7 }t| t
st	 }t&| |&| |'|}|( }z
| )||||}W n t*y- } zt+|  d }~ww tjdd }t"| |d  D ]`\}}|||}| d| } |d urk|,|  }!d urkt-d|!  t.|!||}|rs|dv r{|
/|| || < |d ur|/||  ||kr|
,| }"|0|" t-d|"  |}q?|
j/t1j2|t3d|dd||< |
4|| tjdddks|5  t| ||	S )Nz target must be of GPUTarget typez'source must be either AST or a filepathrF   rI   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_targetr4   r   make_backendr    r5   r   rT   rb   r\   r1   r   r   rS   rJ   r3   rL   rM   rN   rO   r
   r   environgetr   r   r/   	get_groupCompiledKernel__dict__r   
add_stageslistr<   r7   r-   rj   get_codegen_implementationget_module_maprY   	Exceptionr   get_fileprintr   putcreate_location_snapshotr   dumpsvars	put_groupdisable_multithreading)#ri   r   rU   rt   	ir_sourcerT   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_stagerV   rW   rp   r   
use_ir_locr-   
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   )r}   supports_target)r&   r   r   r   r   rA   @  s    z make_backend.<locals>.<listcomp>r   z! compatible backends for target (z) (z). There should only be one.r   )r   valueslenRuntimeErrorrt   )r   activesr   r   r   r   ?  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 rZ   )dataextras)r>   r   r   r   r   r@   I  s   
zLazyDict.__init__r!   r"   c                 C  s0   | j D ]\}}| j|| B | _q| j   | jS rZ   )r   r   clearr>   funcargsr   r   r   r   M  s   
zLazyDict.getc                 C  s   | j ||f d S rZ   )r   r   r  r   r   r   addS  s   zLazyDict.addNr^   )r.   r`   ra   r@   r   r  r   r   r   r   r   G  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)r>   rR   valuer   r   r   __missing__Y  s
   zAsmDict.__missing__N)r.   r`   ra   r
  r   r   r   r   r  W  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   r&   cpr   r   r   r   m  s   $ z*CompiledKernel.__init__.<locals>.<genexpr>cluster_dimsr   rt   arch	warp_sizeKernelMetadatac                 S  s"   g | ]\}}| d st|qS )r   r  r  r   r   r   rA   {  s   " z+CompiledKernel.__init__.<locals>.<listcomp>c                   s:   i | ]}|j d d |j d d  kr| n| qS )r   N)rg   r   rh   )r&   file
binary_extr   r   r)   }  s    ,z+CompiledKernel.__init__.<locals>.<dictcomp>r   )collectionsr  nextr3   r   loadsrh   r8   r   rJ   r   r<   r   r   r   pack_metadatapacked_metadatari   rS   r/   r  r  asmkernelrp   function)r>   ri   r   rS   r  r   r   r   r  rt   	asm_filesr   r  r   r@   k  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drI| jjd urId}| jj|krIt| jj|dtjj| j| j| jj
|\| _ | _| _| _d S )Nmax_shared_memzshared memory	tmem_sizei   ztensor memory)rp   r   r   get_current_devicelauncher_clsri   r   runutilsget_device_propertiessharedr	   rC   r!  load_binaryr/   r  r  n_regsn_spills)r>   device
max_sharedmax_tmem_sizer   r   r   _init_handles  s   

zCompiledKernel._init_handlesc                   s   |dkr|    t |S )Nr$  )r.  super__getattribute__)r>   r/   	__class__r   r   r0    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 ]\}}|| ||< |d7 }q+|| jjj	|| j|f |S )N)r/   r  streamr   r   )r   launch_enter_hookr   r/   r  r4   ri   r    r,   launch_metadatar:   r6   r  r   )	r>   gridr3  r  retarg_dictarg_idxiarg_namer   r   r   r5    s   

zCompiledKernel.launch_metadatac                   s       d d fdd
}|S )N)r3  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_streamr5  r$  r  r  r   r4  launch_exit_hook)r3  r  r+  r5  r6  r>   r   r   runner  s   
"z*CompiledKernel.__getitem__.<locals>.runner)r.  )r>   r6  r?  r   r>  r   __getitem__  s   zCompiledKernel.__getitem__)r.   r`   ra   r4  r=  r@   r.  r0  r5  r@  __classcell__r   r   r1  r   r   d  s    r   )r   r   r]   )/
__future__r   rL   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_patternrk   ptx_arg_type_patternrn   r   r    rb   	lru_cacher   r   r   r   r   r   r1   r  r   r   r   r   r   <module>   sF    
$)
#

&i