a
    khG                    @   s   d dl Z d dlZd dlZd dlZd dlZd dlZd dlZd dlmZ d dl	m
Z
 d dlmZmZmZmZmZmZmZmZmZ ddlmZmZ ddlmZmZ ddlmZmZmZ dd	lmZ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,m-Z-m.Z. dd Z/dd Z0ee1dddZ2ee1dddZ3ee1dddZ4ee1dddZ5ee1dddZ6dd Z7d d! Z8d"d# Z9ee  d$d%d&Z:eej; ee! d'd(d)Z<e1e=e>dhZ?G d*d+ d+Z@G d,d- d-e jAZBG d.d/ d/ZCed0d1G d2d3 d3ZDG d4d5 d5e jAZEd8d6d7ZFdS )9    N)	dataclass)
ModuleType)	AnyCallableDictOptionalTupleTypeUnionIterableList   )knobslanguage)irgluon_ir)	constexpr	str_to_tytensor)_unwrap_if_constexpr
base_value	base_type)get_jit_fn_file_lineget_full_name)JITFunction)find_paths_ifget_iterable_pathset_iterable_path   )CompilationErrorCompileTimeAssertionFailureUnsupportedLanguageConstructc                 C   s*   d}t || s&td| d|  | | S )Nz^[a-zA-Z_][a-zA-Z0-9_]*$zinvalid z identifier: )rematchr   )nametypepattern r'   L/var/www/auris/lib/python3.9/site-packages/triton/compiler/code_generator.pycheck_identifier_legality   s    r)   c                    st   d dd |D }d  fddt D }|dd}|dd}|d	dd
d}|  d| d| }|S )N_c                 S   s   g | ]}|  qS r'   )Zmangle.0tyr'   r'   r(   
<listcomp>!       zmangle_fn.<locals>.<listcomp>c                    s"   g | ]}| d t  |  qS )c)reprr,   i	constantsr'   r(   r.   "   r/   .Z_d_'Z_sq_[]__)joinsortedreplace)r$   Zarg_tysr5   Zmangled_arg_namesZmangled_constantsretr'   r4   r(   	mangle_fn   s    r?   )oreturnc                 C   s
   t | tS N)
isinstancer   r@   r'   r'   r(   _is_triton_value+   s    rE   c                 C   s
   t | tS rB   )rC   r   rD   r'   r'   r(   _is_triton_tensor/   s    rF   c                 C   s   | d u pt | ttjjtfS rB   )rC   r   r   coredtyper   rD   r'   r'   r(   _is_constexpr3   s    rI   c                 C   s   t | o| j o| jjdkS Nr   )rF   r%   is_blockZnumelrD   r'   r'   r(   _is_non_scalar_tensor7   s    rL   c                 C   s   t | ttfS rB   )rC   listtuplerD   r'   r'   r(   _is_list_like;   s    rO   c              
   C   sT   |j rPt|D ]@\}}t|st|rt|j| d|j d|j|  d| qd S )Nz	Function z= is marked noinline, but was called with non-scalar argument :)noinline	enumeraterI   rL   r!   src__name__	arg_names)nodefnargsidxargr'   r'   r(   _check_fn_args?   s    r[   c                 C   s   t | tot| tot| dS )N_fields)rC   r%   
issubclassrN   hasattr)valr'   r'   r(   _is_namedtupleI   s    r`   c                    s   t t| r| j}n,t| tjr*| jj}nds@J dt|   fdd| D }dd |D }dd |D }t|t||S )NFzUnsupported type c                    s   g | ]} |qS r'   r'   r,   vrW   r'   r(   r.   U   r/   z*_apply_to_tuple_values.<locals>.<listcomp>c                 S   s    g | ]}|d u rt |n|qS rB   )r   ra   r'   r'   r(   r.   V   r/   c                 S   s   g | ]
}|j qS r'   r%   ra   r'   r'   r(   r.   W   r/   )r`   r%   r\   rC   r   rN   fields
tuple_type)valuerW   re   valstypesr'   rc   r(   _apply_to_tuple_valuesM   s    
rj   )valuesc                 C   s   g }| D ]}| | q|S rB   )_flatten_ir)rk   handlesrb   r'   r'   r(   flatten_values_to_ir[   s    rn   )rm   ri   c                 c   s8   d}|D ]}| | |\}}|V  q|t| ks4J d S )Nr   )_unflatten_irlen)rm   ri   cursorr-   rg   r'   r'   r(   unflatten_ir_valuesb   s
    rr   c                   @   s$   e Zd Zdd Zdd Zdd ZdS )enter_sub_regionc                 C   s
   || _ d S rB   )	generator)selfrt   r'   r'   r(   __init__o   s    zenter_sub_region.__init__c                 C   sL   | j j | _| j j | _i | j _| j j | _| j j	 | _
| j| jfS rB   )rt   lscopecopyliveins
local_defs	prev_defsbuilderget_insertion_blockinsert_blockget_insertion_pointinsert_pointru   r'   r'   r(   	__enter__r   s    zenter_sub_region.__enter__c                 O   s(   | j j| j | j| j _| j| j _d S rB   )rt   r|   restore_insertion_pointr   ry   rw   r{   rz   )ru   rX   kwargsr'   r'   r(   __exit__{   s    
zenter_sub_region.__exit__N)rT   
__module____qualname__rv   r   r   r'   r'   r'   r(   rs   m   s   	rs   c                   @   s   e Zd Zdd ZedddZedddZeddd	Zej	ed
ddZ
ejed
ddZejed
ddZejed
ddZejed
ddZejed
ddZejed
ddZejed
ddZejed
ddZejed
ddZdS ) ContainsReturnCheckerc                 C   s
   || _ d S rB   )gscope)ru   r   r'   r'   r(   rv      s    zContainsReturnChecker.__init__rA   c                    s   t  fdd|D S )Nc                 3   s   | ]}  |V  qd S rB   visit)r,   sr   r'   r(   	<genexpr>   r/   z5ContainsReturnChecker._visit_stmts.<locals>.<genexpr>)any)ru   bodyr'   r   r(   _visit_stmts   s    z"ContainsReturnChecker._visit_stmtsc                 C   s   dS NFr'   )ru   rW   r'   r'   r(   _visit_function   s    z%ContainsReturnChecker._visit_functionc                 C   sf   d}t |D ]R\}}t|trF|D ]}t|t jr$|p@| |}q$qt|t jr|p^| |}q|S r   )astiter_fieldsrC   rM   ASTr   )ru   rV   r>   r*   rg   itemr'   r'   r(   generic_visit   s    
z#ContainsReturnChecker.generic_visitrV   rA   c                 C   sP   t |jtjrD|jj| jv r@| j|jj }t||j}| |S dS | 	|jS r   )
rC   rg   r   Nameidr   getattrattrr   r   )ru   rV   rg   rW   r'   r'   r(   visit_Attribute   s    
z%ContainsReturnChecker.visit_Attributec                 C   s:   t |jtju rdS |j| jv r6| j|j }| |S dS r   )r%   ctxr   Storer   r   r   )ru   rV   rW   r'   r'   r(   
visit_Name   s    
z ContainsReturnChecker.visit_Namec                 C   s   dS NTr'   ru   rV   r'   r'   r(   visit_Return   s    z"ContainsReturnChecker.visit_Returnc                 C   s   dS r   r'   r   r'   r'   r(   visit_Assign   s    z"ContainsReturnChecker.visit_Assignc                 C   s   dS r   r'   r   r'   r'   r(   visit_AugAssign   s    z%ContainsReturnChecker.visit_AugAssignc                 C   s   |  |jS rB   r   r   r   r'   r'   r(   visit_Module   s    z"ContainsReturnChecker.visit_Modulec                 C   s   |  |jS rB   r   r   r'   r'   r(   visit_FunctionDef   s    z'ContainsReturnChecker.visit_FunctionDefc                 C   s&   |  |j}|jr"|p |  |j}|S rB   )r   r   orelse)ru   rV   r>   r'   r'   r(   visit_If   s    zContainsReturnChecker.visit_Ifc                 C   s   |  |jp|  |jS rB   )r   r   r   r   r'   r'   r(   visit_IfExp   s    z!ContainsReturnChecker.visit_IfExpc                 C   s   |  |jS rB   )r   funcr   r'   r'   r(   
visit_Call   s    z ContainsReturnChecker.visit_CallN)rT   r   r   rv   boolr   r   r   r   	Attributer   r   r   Returnr   Assignr   	AugAssignr   Moduler   FunctionDefr   Ifr   IfExpr   Callr   r'   r'   r'   r(   r      s   r   c                   @   sb   e Zd Zdd Zejee eej dddZ	ejeej dddZ
ejd	d
dZdd ZdS )ASTFunctionc                 C   s   || _ || _|| _|| _d S rB   )	ret_types	arg_typesr5   attrs)ru   r   r   r5   r   r'   r'   r(   rv      s    zASTFunction.__init__)r|   ri   rA   c                 C   s(   g }|D ]}|d u rq| || q|S rB   )Z_flatten_ir_types)ru   r|   ri   Zir_typesr-   r'   r'   r(   flatten_ir_types   s    zASTFunction.flatten_ir_types)r|   rA   c                 C   s   |  || jS rB   )r   r   )ru   r|   r'   r'   r(   return_types_ir   s    zASTFunction.return_types_ir)r|   c                    sP    fdd}t t j|} fdd|D } ||} |}|||S )Nc                    s   |  j vo|d uS rB   r4   pathr*   r   r'   r(   <lambda>   r/   z'ASTFunction.serialize.<locals>.<lambda>c                    s   g | ]}t  j|qS r'   )r   r   r,   r   r   r'   r(   r.      r/   z)ASTFunction.serialize.<locals>.<listcomp>)rM   r   r   r   r   Zget_function_ty)ru   r|   is_val	val_pathsr   Zarg_types_irZret_types_irr'   r   r(   	serialize   s    
zASTFunction.serializec                    s   fddj }fdd}ttj |}d} fddt  D }|D ]V}tj |}j|g }	|	D ]\}
} ||
| qv|	||\}}t
||| qTj}| D ]\}}t
||t| q|S )Nc                    s6   t | tttjfr,t fdd| D | S td S )Nc                    s   g | ]} |qS r'   r'   r,   xmake_templater'   r(   r.      r/   zBASTFunction.deserialize.<locals>.make_template.<locals>.<listcomp>)rC   rM   rN   r   rf   r   )r-   r   r'   r(   r      s    z.ASTFunction.deserialize.<locals>.make_templatec                    s   |  j vo|d uS rB   r4   r   r   r'   r(   r      r/   z)ASTFunction.deserialize.<locals>.<lambda>r   c                    s   g | ]}  |qS r'   rX   r2   rc   r'   r(   r.     r/   z+ASTFunction.deserialize.<locals>.<listcomp>)r   rM   r   rangeZget_num_argsr   r   getZset_arg_attrro   r   r5   itemsr   r   )ru   rW   rh   r   r   rq   rm   r   r-   Z
attr_specs	attr_nameZattr_valr_   r5   r'   )rW   r   ru   r(   deserialize   s"    
zASTFunction.deserializeN)rT   r   r   rv   r   r|   r   r   r%   r   r   r   r   r'   r'   r'   r(   r      s
   
r   T)frozenc                   @   s   e Zd ZU eed< eed< dS )BoundJITMethod__self____func__N)rT   r   r   r   __annotations__r   r'   r'   r'   r(   r     s   
r   c                       s  e Zd ZU deee ee dddZdd ee	e
eeeeefD Zeeef ed	< ed
ejjfdejfdejff dd Zdd Zdd Zeeeef ddddZdd Z dd Z!dd Z"dd Z#dd Z$d d! Z%ed"d#d$Z&d%d& Z'd'd( Z(d)d* Z)d+d, Z*d-d. Z+d/d0 Z,d1d2 Z-d3d4 Z.d5d6 Z/d7d8 Z0d9d: Z1d;d< Z2d=d> Z3e4j5d?e4j6d@e4j7dAe4j8dBe4j9dCe4j:dDe4j;dEe4j<dFe4j=dGe4j>dHe4j?dIe4j@dJiZAeeBe4jC ef edK< dLdM ZDdNdO ZEdPdQ ZFdRdS ZGdTdU ZHdVdW ZIdXdY ZJe4jKdZe4jLd[e4jMd\e4jNd]e4jOd^e4jPd_iZQeeBe4jR ef ed`< dadb ZSe4jTdce4jUdde4jVdee4jWdfiZXeeBe4jY ef edg< dhdi ZZdjdk Z[dldm Z\dndo Z]dpdq Z^drds Z_dtdu Z`dvdw Zadxdy Zbdzd{ Zcd|d} Zdeeeef d"d~dZfed"ddZgedddZhdd Zidd Zje4jkdddZle4jmde4jndiZoeeBe4jp ef ed< dd Zqdd Zrdd Zsdd Zt fddZudd Zve4jwddddZxdd ZyddlzmZ{ ejj|exejj}eye~e{j|exe{j}eye~eeyeeeyeiZeeee4jwgef f ed<   ZS )CodeGeneratorNFr   )jit_fnfunction_types	file_namec                 C   s  || _ | r4ddlm} t|| _|| j| _n$ddlm	} t
|| _|| j| _|| _|d | _| j||d || j_|| j_|d u ri n|| j_|	d u r| j n|	| _|d u ri n|| _|| _i | _| D ]b\}}t|tr||j|| j|< qt|dd}||v r2t|| |j| j|< q|| j|< qi | _|| _|
rp||dd d  }t|d}|| _ |
| _!d | _"|| _#g | _$d | _%i | _&| ' | _(d | _)d	| _*d S )
Nr   )GluonSemantic)TritonSemanticr   r    r6   functionF)+contextZis_gluonZ,triton.experimental.gluon.language._semanticr   r   ZGluonOpBuilderr|   semanticZtriton.language.semanticr   r   r   
begin_lineset_locoptionscodegen_fns
module_mapcreate_modulemodulefunction_ret_types	prototyper   r   rC   r   r   rT   r   rw   r   rfindr)   function_name	is_kernelcur_noderQ   	scf_stackret_typerz   _define_name_lookupdereference_namerW   visiting_arg_default_value)ru   r   r   r   r   r   r   r   r   r   r   r   rQ   r   r   r   r   krb   module_namer'   r'   r(   rv     sR    



zCodeGenerator.__init__c                 C   s   i | ]}|j |qS r'   rT   r,   r*   r'   r'   r(   
<dictcomp>V  s   zCodeGenerator.<dictcomp>builtin_namespaceprintminmaxc                 C   s   t | jj||S rB   )r!   r   rS   )ru   rV   messager'   r'   r(   _unsupported`  s    zCodeGenerator._unsupportedc                 C   s0   t  }| j||}||u r dS t|r,dS dS )NFT)objectr   r   rI   )ru   r$   absent_markerr_   r'   r'   r(   _is_constexpr_globalc  s    z"CodeGenerator._is_constexpr_globalc                    sH   t dfddt dfddt  t td fdd}|S )	N)r$   c                    s    j | |S rB   )rw   r   )r$   absentr   r'   r(   local_lookupp  s    z7CodeGenerator._define_name_lookup.<locals>.local_lookupc                    s    j | |}t||u |  jv t|tu t|tt|ddt|ddt|dd	dt|dd	dt|t
jt| |  jtjjgr|S ttd|  d	d
dd S )NZ__triton_builtin__FZ__triton_aggregate__r   r   ztriton.languagez"triton.experimental.gluon.languagez.                Cannot access global variable a   from within @jit'ed
                function. Triton kernels can only access global variables that
                are instanstiated as constexpr (`x = triton.language.constexpr(42)`). Note that this is different from
                annotating a variable as constexpr (`x: triton.language.constexpr = 42`), which is not supported.  Alternatively, set the
                envvar TRITON_ALLOW_NON_CONSTEXPR_GLOBALS=1, but we do not
                promise to support this forever.
 )r   r   r   r   r%   r   rC   r   r   
startswithr   rH   r`   r   r   r   compilationZallow_non_constexpr_globals	NameErrortextwrapdedentr=   )r$   r   r_   r   r'   r(   global_lookupt  s0    




z8CodeGenerator._define_name_lookup.<locals>.global_lookup)r$   rA   c                    sD    }j jfD ]}|| |}||ur|  S qt|  dd S )Nz is not defined)r   r   r  )r$   r   Zlookup_functionrg   r   r  r   ru   r'   r(   name_lookup  s    

z6CodeGenerator._define_name_lookup.<locals>.name_lookup)strr   r   )ru   r
  r'   r	  r(   r   n  s
    !z!CodeGenerator._define_name_lookup)r$   rg   rA   c                 C   s   || j |< || j|< dS )z This function:
            called by visit_Assign() & visit_FunctionDef() to store left value (lvalue)
        1. record local defined name (FIXME: should consider control flow)
        2. store tensor in self.lvalue
        N)rw   rz   )ru   r$   rg   r'   r'   r(   	set_value  s    
zCodeGenerator.set_valuec                 C   s   | j  }| j  }||fS rB   )r|   get_locr   )ru   locipr'   r'   r(   _get_insertion_point_and_loc  s    

z*CodeGenerator._get_insertion_point_and_locc                 C   s   | j | | j | d S rB   )r|   r   r   )ru   r  r  r'   r'   r(   _set_insertion_point_and_loc  s    z*CodeGenerator._set_insertion_point_and_locc                 C   s6   t |s|g}|D ]}| | t|tjr q2qd S rB   )rO   r   rC   r   r   )ru   Zstmtsstmtr'   r'   r(   visit_compound_statement  s    
z&CodeGenerator.visit_compound_statementc                 C   s   t j| | d S rB   r   NodeVisitorr   r   r'   r'   r(   r     s    zCodeGenerator.visit_Modulec                    s6     |j}|d u sJ t fdd|jD }|S )Nc                    s   g | ]}  |qS r'   r   )r,   eltr   r'   r(   r.     r/   z,CodeGenerator.visit_List.<locals>.<listcomp>)r   r   r   rN   elts)ru   rV   r   r  r'   r   r(   
visit_List  s    zCodeGenerator.visit_Listc                    s    |j}g } fdd  |}|d u r6tj}n"t|tjjsHJ || |j}j	
| jd u rv|_n j|krtdj d| j	 }j	| d S )Nc                    s8   t | tjrt|  S t | tjttfr4j| S | S rB   )	rC   r   rN   rj   r   intfloatr   	to_tensor)rg   decayru   r'   r(   r    s
    
z)CodeGenerator.visit_Return.<locals>.decayzInconsistent return types:  and )r   rg   r   voidrC   rG   r   rl   r%   r|   r>   r   	TypeErrorcreate_blockset_insertion_point_to_end)ru   rV   Z	ret_valuerm   Zret_tyZpost_ret_blockr'   r  r(   r     s     



zCodeGenerator.visit_Returnr   c                 C   s$   |  |j}t|tjjsJ |jS rB   )r   rg   rC   r   rG   rN   rk   ru   rV   rX   r'   r'   r(   visit_Starred  s    zCodeGenerator.visit_Starredc              	      s4    |j\}} jr" |dt|jjd d d D ]\}}|jj| d  }|j}|j}tj	|t
 d}	|d u rtj|	g|d}
ntj|	||d}
z$ jrJ d _  |
 W d _q8d _0 q8 jrd	nd
} j j} j j j|| j _ j j  j } j j}t||D ]\}} || q2 j } j|  |j  j  rJ  j d u s j t!j"krt!j" _  j#g  ndt$ j t!j%rΈ j j& j_'n j g j_' j( j j  j# fdd j) jD   j*  |r0 j+| d S )Nz,nested function definition is not supported.r   )r   r   targetsrg   )targetrg   
annotationTFpublicprivatec                    s   g | ]} j |qS r'   )r|   create_poisonr+   r   r'   r(   r.      r/   z3CodeGenerator.visit_FunctionDef.<locals>.<listcomp>),r   rX   rW   r   rR   defaultsr)  rZ   r   r   r   r   	AnnAssignr   r   r   r   r|   Zget_or_insert_functionr   r   rQ   Z	push_backZadd_entry_blockr   zipr  r}   set_insertion_point_to_startr  r   has_terminatorr   r   r  r>   rC   rf   ri   r   Z
reset_typer   finalizer"  )ru   rV   rU   kwarg_namesr3   default_valueZarg_noder)  r$   Z	st_targetZ	init_nodeZ
visibilityZfn_tyentryZ
arg_valuesZarg_name	arg_valueZ	insert_ptr'   r   r(   r     sN    


$
zCodeGenerator.visit_FunctionDefc                 C   s4   g }|j D ]}|| |g7 }q
| |j}||fS rB   )rX   r   kwarg)ru   rV   rU   rZ   r3  r'   r'   r(   visit_arguments&  s
    
zCodeGenerator.visit_argumentsc                 C   s   t j| | |jS rB   )r   r  r   rZ   r   r'   r'   r(   	visit_arg-  s    zCodeGenerator.visit_argc                 C   sj   |  |j}|  |j}|  |j}|tkr`|| jv rDt| dt|}|| j|< | j| S | |S )Nz4 is already defined. constexpr cannot be reassigned.)r   r)  r(  rg   r   rw   
ValueErrorr   )ru   rV   r)  r(  rg   r'   r'   r(   visit_AnnAssign1  s    


zCodeGenerator.visit_AnnAssignc                 C   s   t |jtjsJ t |tjr*| ||S t |tjr`t|jD ]\}}| 	||j
|  q@d S t |tjr| |j}t||j| d S t |tjsJ | | || d S rB   )rC   r   r   r   	Subscriptvisit_Subscript_Storer   rR   r  assignTargetrk   r   r   rg   setattrr   r   r  )ru   r(  rg   r3   baser'   r'   r(   r>  A  s    zCodeGenerator.assignTargetc                    s\    fdd   |j}t|tjr2|jgn|j}t|dksHJ |d | d S )Nc                    sT   t | tjrt|  S tjtjf}t| } | d urPt| sPt | |sPj| } | S rB   )	rC   r   rN   rj   rH   r   rE   r   r  )rg   Znative_nontensor_types_sanitize_valueru   r'   r(   rB  R  s    
z3CodeGenerator.visit_Assign.<locals>._sanitize_valuer   r   )	r   rg   rC   r   r.  r(  r'  rp   r>  )ru   rV   rk   r'  r'   rA  r(   r   P  s
    zCodeGenerator.visit_Assignc                 C   sN   t |j}t |_t||j|j}tj	|jg|d}| 
| | 
|S )Nr&  )rx   deepcopyr(  r   Loadr   BinOpoprg   r   r   )ru   rV   lhsrhsZassignr'   r'   r(   r   b  s    

zCodeGenerator.visit_AugAssignc                 C   s"   t |jtju r|jS | |jS rB   )r%   r   r   r   r   r   r   r'   r'   r(   r   j  s    zCodeGenerator.visit_Namec                 C   s   t j| | d S rB   r  r   r'   r'   r(   visit_Storeo  s    zCodeGenerator.visit_Storec                 C   s   t j| | d S rB   r  r   r'   r'   r(   
visit_Loadr  s    zCodeGenerator.visit_Loadc                    s    fdd|j D }t|S )Nc                    s   g | ]}  |qS r'   r   r   r   r'   r(   r.   v  r/   z-CodeGenerator.visit_Tuple.<locals>.<listcomp>)r  r   rN   r#  r'   r   r(   visit_Tupleu  s    zCodeGenerator.visit_Tuplec                 C   sv   t |rt|||| jdS t |rFtdd|}t|||| jdS t|ttjfsht|trht|}t|||S )N	_semanticz__(.*)__z__r\1__)	rF   r   r   r"   subrC   r   r   rN   )ru   method_namerG  rH  Zreverse_method_namer'   r'   r(   _apply_binary_methody  s    z"CodeGenerator._apply_binary_methodc                 C   sV   |  |j}|  |j}| jt|j}|d u rH| |d|jj	| 
|||S )Nz8AST binary operator '{}' is not (currently) implemented.)r   leftright_method_name_for_bin_opr   r%   rF  r   formatrT   rP  )ru   rV   rG  rH  rO  r'   r'   r(   visit_BinOp  s    zCodeGenerator.visit_BinOp__add____sub____mul____truediv____floordiv____mod____pow__
__lshift__
__rshift____and____or____xor__rS  c                 C   s  | j | | |j | j  }| j }i }|jrr| j | | | _i | _| |j | j }| j  }g }|D ]}|df|dffD ]h\}	}
||	v rt	|	| t	|| k}|r|	| j	|| j	ksJ d| d||  d|
 d|	|  q||v s||v r|
| ||v r6||vr6|| ||< ||v rz||vrz|| ||< qzt| | @ D ]v}||v rzqh|| }|j	}|| }|j	}t	|t	|k}|r||ksJ d| d| d	| d
|
| qh|||||fS )NZthenelsezinitial value for `z` is of type z
, but the z block redefines it as zMismatched type for z between then block (z) and else block ())r|   r0  r  r   r}   rz   rx   r   rw   r%   appendr<   keys)ru   rV   ry   
then_block
else_block	then_defs	else_defsnamesr$   ZdefsZ
block_nameZ
type_equalthen_valZthen_tyelse_valZelse_tyr'   r'   r(   visit_then_else_blocks  sZ    






z$CodeGenerator.visit_then_else_blocksc                    s  t | B}|\}}| j }| j }| j| | j|j|| | ||||\ }}}| j | j| | rJ | tfdd|D }	| j	|	 | j| | rJ | t fdd|D }
| j	|
 t
|	t
|
ksJ t|	|
D ].\}}| }|| ks0J | qW d    n1 sT0    Y  | j fddtt
|	D }fdd|D }t||}t||D ]\}}| || qd S )Nc                 3   s   | ]} | V  qd S rB   r'   r,   r$   rh  r'   r(   r     r/   z3CodeGenerator.visit_if_top_level.<locals>.<genexpr>c                 3   s   | ]} | V  qd S rB   r'   rn  ri  r'   r(   r     r/   c                    s   g | ]}  |qS r'   rZ   r2   )endif_blockr'   r(   r.     r/   z4CodeGenerator.visit_if_top_level.<locals>.<listcomp>c                    s   g | ]} | j qS r'   rd   rn  ro  r'   r(   r.     r/   )rs   r|   r!  r"  Zcreate_cond_branchhandlerm  r1  rn   Zcreate_branchrp   r/  get_typeadd_argumentr0  r   rr   r  )ru   condrV   srry   Zip_blockrf  rg  rj  then_handleselse_handlesZthen_hZelse_hr-   res_handlesri   
new_valuesr$   	new_valuer'   )ri  rr  rh  r(   visit_if_top_level  s8    


.
z CodeGenerator.visit_if_top_levelc                    s  t | .}|\}}|  \}}| j }|jr:| j nd }	| ||||	\ }}	}
tfdd|
D }| || | jdd |D |j	d|
  | j  t|
dkr| j| |js܈ }	n|	
  | j  t|
dkr*t fdd|
D }| j| W d    n1 s@0    Y  fddtt|D }fd	d|
D }t||}t|
|D ]\}}| || qd S )
Nc                 3   s   | ]} | V  qd S rB   r'   rn  ro  r'   r(   r      r/   z-CodeGenerator.visit_if_scf.<locals>.<genexpr>c                 S   s   g | ]}|  qS r'   rt  r,   hr'   r'   r(   r.     r/   z.CodeGenerator.visit_if_scf.<locals>.<listcomp>Tr   c                 3   s   | ]} | V  qd S rB   r'   rn  rp  r'   r(   r     r/   c                    s   g | ]}  |qS r'   
get_resultr2   )if_opr'   r(   r.     r/   c                    s   g | ]} | j qS r'   rd   rn  ro  r'   r(   r.     r/   )rs   r  r|   r!  r   rm  rn   r  create_if_oprs  merge_block_beforeget_then_blockr"  rp   create_yield_opget_else_blockr   rr   r/  r  )ru   rv  rV   rw  ry   r*   r  last_locrf  rg  rj  rx  ry  rz  ri   r{  r$   r|  r'   )ri  r  rh  r(   visit_if_scf  s4    

,
zCodeGenerator.visit_if_scfc              	   C   s  |  |j}t|rt|r(| |d|j r\tdt	
|j  tjj|| j| d}|jtj| jd}t| j |}|r| jr| |d| || n| || nXt|}t|tvr| |dddd	 tD t|j|r|jn|j}| | d S )
Nz=Boolean value of Tensor with more than one value is ambiguousziIf conditional called with multidimensional Tensor instead of scalar; please use "if (%s).item()" instead)rM  
_generatorrL  zCannot have `return` statements inside `while` or `for` statements in triton (note that this also applies to `return` statements that are inside functions transitively called from within `while`/`for` statements)O`if` conditionals can only accept values of type {{{}}}, not objects of type {}, c                 s   s   | ]}|j V  qd S rB   r   r   r'   r'   r(   r   2  r/   z)CodeGenerator.visit_If.<locals>.<genexpr>)r   testrF   rL   r   r%   rK   warningswarnr   unparser   rG   Z_unsplatr   toint1r   r   r   r}  r  r   _condition_typesrT  r;   rT   r   r   r  )ru   rV   rv  Zcontains_returnZactive_blockr'   r'   r(   r     s<    

zCodeGenerator.visit_Ifc              	   C   s&  |  |j}t|r|jtj| jd}t| ~ |  \}}| j	
 }| j	| | j|  |j}| j	 }| j	
 }| j	| | j|  |j}| j	 }| || |j|jksJ d|j d|j |j}	|	tjkr|	| j	gng }
| j	|
|jd}||  |
r@| j	|  | j	|jg | j	|  ||  |
r| j	|  | j	|jg |
rtj|d|	nd W  d    S 1 s0    Y  n^t|}t|tvr|  |d!d"dd	 tD t|j#|r|  |jS |  |jS d S )
NrL  zATernary expression with dynamic condition has inconsistent types r  Tr   r  r  c                 s   s   | ]}|j V  qd S rB   r   r   r'   r'   r(   r   f  r/   z,CodeGenerator.visit_IfExp.<locals>.<genexpr>)$r   r  rF   r  r   r  r   rs   r  r|   r!  r0  r  r   r}   r   r  r%   r  to_irr  rs  r  r  r"  r  r  rG   r   r  r   r  r   rT  r;   rT   )ru   rV   rv  r  r  rf  rk  rg  rl  r   Zret_type_irr  r'   r'   r(   r   8  sR    




@zCodeGenerator.visit_IfExpc                 C   s   d S rB   r'   r   r'   r'   r(   
visit_Passm  s    zCodeGenerator.visit_Passc                 C   s   t |jdkrt |jdks(| |d| |j}| |jd }t|}t|}t|jd tj	u rtt
||u S t|jd tju rt
||uS | jt|jd }|d u r| |d|jd j| |||S )Nr   z1simultaneous multiple comparison is not supportedr   z<AST comparison operator '{}' is not (currently) implemented.)rp   comparatorsopsr   r   rQ  r   r%   r   Isr   IsNot_method_name_for_comp_opr   rT  rT   rP  )ru   rV   rG  rH  	lhs_value	rhs_valuerO  r'   r'   r(   visit_Comparep  s     zCodeGenerator.visit_Compare__eq____ne____lt____le____gt____ge__r  c                 C   s   |  |j}| jt|j}|d u r>| |d|jj dt|rXt	||| j
dS zt	|| W S  ty   |dkrt|  Y S | |d| dt|j Y n0 d S )NzAST unary operator 'z!' is not (currently) implemented.rL  __not__z)' is not (currently) implemented on type )r   operand_method_name_for_unary_opr   r%   rF  r   rT   rF   r   r   AttributeErrorr   )ru   rV   r  rW   r'   r'   r(   visit_UnaryOp  s    zCodeGenerator.visit_UnaryOp__neg____pos__r  
__invert__r  c                 C   s   t |sJ d| dt |s0J d| dt|t|u sPJ d| dt|r|j|jksJ d| d|j d|j d	d S )
Nzcannot reassign constxpr z in the loopzcannot reasign constexpr zLoop carried variable z changed typezLoop-carried variable z has initial type z but is re-assigned to z: in loop! Please make sure that the type stays consistent.)rE   r%   rF   )ru   r$   loop_vallive_valr'   r'   r(   _verify_loop_carried_variable  s     z+CodeGenerator._verify_loop_carried_variablec                 C   s   |  |jS rB   )r   context_exprr   r'   r'   r(   visit_withitem  s    zCodeGenerator.visit_withitemc                    s   t |jdksJ |jd j} |j}|tjkr fdd|jD }||d ji  	|j
 W d    q1 sz0    Y  n 	|j
 d S )Nr   r   c                    s   g | ]}  |qS r'   r   r,   rZ   r   r'   r(   r.     r/   z,CodeGenerator.visit_With.<locals>.<listcomp>Z_builder)rp   r   r  r   r   r   Z
async_taskrX   r|   r  r   )ru   rV   r   ZwithitemClassrX   r'   r   r(   
visit_With  s    
,zCodeGenerator.visit_Withc                    s  t | P}|\}}|  \}}| j }| j| | j| | |j | j	  | j
}|  g }	g }
|D ]>}||v rr|| }|| }| ||| |	| |
| qrt|
}dd |D }dd |
D }| || | j||| j || j fddtt|D }t||}t|	|D ]\}}|| j|< || j
|< q<| |j}| j | j|j| | j | | j   fddtt|D }t||}t|	|D ]\}}|| j|< || j
|< q| j| | |j | j	  | j
}g }|D ]}||v r || | q | j| W d    n1 sb0    Y  fddtt|D }t||}t|	|D ]\}}|| j|< || j
|< q|jD ]"}dsJ dt j!"| | qd S )	Nc                 S   s   g | ]}|  qS r'   r~  r  r'   r'   r(   r.     r/   z-CodeGenerator.visit_While.<locals>.<listcomp>c                 S   s   g | ]
}|j qS r'   rd   )r,   ar'   r'   r(   r.     r/   c                    s   g | ]}  |qS r'   rq  r2   )before_blockr'   r(   r.     r/   c                    s   g | ]}  |qS r'   rq  r2   )after_blockr'   r(   r.     r/   c                    s   g | ]}  |qS r'   r  r2   )while_opr'   r(   r.     r/   FzNot implemented)#rs   r  r|   r!  r0  r   rd  r  r   poprz   eraser  rn   r  Zcreate_while_opZcreate_block_with_parentZ
get_beforer   rp   rr   r/  rw   r   r  r"  Zcreate_condition_oprs  Z	get_afterrl   r  r   r   r  r   )ru   rV   rw  ry   r~   r  r  dummyZ	loop_defsrj  	init_argsr$   r  r  init_handlesinit_tysZinit_fe_tys
block_argsZcondition_argsr_   rv  Zbody_handlesZ	body_argsyieldsresult_handlesZresult_valsZnew_defr  r'   )r  r  r  r(   visit_While  st    








,


zCodeGenerator.visit_Whilec                 C   sJ   t |jtjsJ | |j}| |j}t|rB|j|| j	dS || S )NrL  )
rC   r   r   rD  r   rg   slicerF   __getitem__r   )ru   rV   rG  slicesr'   r'   r(   visit_Subscript_Load  s    z"CodeGenerator.visit_Subscript_Loadc                 C   sJ   t |jtjsJ | |j}| |j}t |tjs:J |	|| d S rB   )
rC   r   r   r   r   rg   r  r   rN   __setitem__)ru   rV   rg   rG  r  r'   r'   r(   r=    s
    z#CodeGenerator.visit_Subscript_Storec                 C   s
   |  |S rB   )r  r   r'   r'   r(   visit_Subscript  s    zCodeGenerator.visit_Subscriptc                    s    fdd|j D S )Nc                    s   g | ]}  |qS r'   r   )r,   Zdimr   r'   r(   r.     r/   z0CodeGenerator.visit_ExtSlice.<locals>.<listcomp>)dimsr   r'   r   r(   visit_ExtSlice  s    zCodeGenerator.visit_ExtSlicec           ,         st   |jj}fdd|jjD }tfdd|jjD }|tjkr||i |}t|j	j
|jj
|jj
}|D ]<}t|j|jj< |j |jD ]}tj| qqrd S d }	d }
d}d}d}|tju r||i |}|j	}|j}|j}|j}	|j}
|j}|j}|j}n|tu rt|dkr4|d n td}t|dkrZ|d n |jjd }t|dkr|d n td}ntd	d}t|r|j
dk rt|j
 }d
}|| }}j !|}j !|}j !|}|j"# r|j"# r|j"# s:t$d|j" d|j" d|j" dj %|j"|j"}j %||j"}|&j'}|j(tj)j"j*j+k}|j,}|j,}|j,}j'-|||}j'-|||}j'-|||}j'.|}/|jjtj)0|| t1}|\}}2 \}}j'3 }j'4| j56| |j j57  |8  g }g }g }j9D ]N}||v rTj9| } || }!:|| |! |6| |6|! |6|  qT;|| t<|}"dd |D }#j'=||||" t>|	d ur ?dj'@|	 t>|
d ur ?dj'@|
 |r4 ?dj'A  |rL ?dj'A  |rd ?dj'A  j56|  Bdj'4 |C _i _9fddtt|"D }$tD|$|#}%tE||%D ]\}}&/||& qĈ|j j57  g }j9D ]<}||v rj9| }'tF|'tr,j !|'}'|6|' qt|dkr\t<|}(j'G|( H })|)I dkszJ dj'4  J }|rj'K||}j'L||}j|jj j,M| /|jjtj)0|| W d    n1 s0    Y   fddtt|"D }*tD|*|#}+tE||+D ]\}}&/||& q.|jD ]"}ds^J dtj| qLd S )Nc                    s   g | ]}  |qS r'   r   r  r   r'   r(   r.     r/   z+CodeGenerator.visit_For.<locals>.<listcomp>c                 3   s   | ]}  |V  qd S rB   r   r,   keywordr   r'   r(   r     r/   z*CodeGenerator.visit_For.<locals>.<genexpr>Fr   r   r   zAOnly `range` and `static_range` iterators are currently supportedTz0For loop bounds and step must all be ints, are (r  rc  c                 S   s   g | ]
}|j qS r'   rd   ra   r'   r'   r(   r.   u  r/   ztt.num_stagesztt.loop_unroll_factorztt.disallow_acc_multi_bufferz
tt.flattenztt.warp_specializec                    s   g | ]}  |d  qS )r   rq  r2   )for_op_bodyr'   r(   r.     r/   z7We use SCF, so the loop body should only have one blockc                    s   g | ]}  |qS r'   r  r2   )for_opr'   r(   r.     r/   z)Don't know what to do with else after for)Nr   iterr   rX   dictkeywordsr   static_ranger   startrg   endstepr   rw   r(  r   r  r   r   r   r  r   
num_stagesloop_unroll_factordisallow_acc_multi_bufferflattenwarp_specializerp   NumRuntimeErrorrI   r   r  rH   Zis_intr   Zinteger_promote_implr  r|   Zint_signednessrG   Z
SIGNEDNESSZSIGNEDrs  Zcreate_int_castr,  r  r   rs   r  r!  r0  r   rd  r  r  rz   r  r  rn   Zcreate_for_opr   Zset_attrZget_int32_attrZget_unit_attrget_bodyrx   rr   r/  rC   r  Z
get_parentsizeZget_induction_varZ
create_subZ
create_addZreplace_all_uses_with),ru   rV   ZIteratorClassZ	iter_argsZiter_kwargsiteratorr  r3   r  r  r  r  r  r  ZlbZubr  Znegative_stepZiv_typeZ
iv_ir_typeZiv_is_signedZivrw  ry   r~   r  r  blockr  r  rj  r$   r  r  r  r  Zblock_handlesr  r_   localZyield_handlesZfor_op_regionr  Zresult_valuesr'   )r  r  ru   r(   	visit_For  s    


&((
$"













:

zCodeGenerator.visit_Forc                 C   s2   |  |j}|  |j}|  |j}t|||S rB   )r   lowerupperr  r   r  )ru   rV   r  r  r  r'   r'   r(   visit_Slice  s    zCodeGenerator.visit_Slicec                 C   s   |  |jS rB   )r   rg   r   r'   r'   r(   visit_Index  s    zCodeGenerator.visit_Indexc                 C   s   |j | |jfS rB   )rZ   r   rg   r   r'   r'   r(   visit_keyword  s    zCodeGenerator.visit_keywordc                 C   s:   |  |j}|jd ur"|  |jnd}tjj||| jdS )Nr   rL  )r   r  msgr   rG   Zdevice_assertr   )ru   rV   r  r  r'   r'   r(   visit_Assert  s    zCodeGenerator.visit_Assertrc   c                    s  t j|jg R i |  fdd|jD  t D ].\}}t|tjtt	t
tfr6tj| |< q6t dd } fdd|D }t dd } fdd|D }tt|d	d |D |}	| j|	st|\}
}d
d  D }tg ||t }t| j|| | j||	| j|j|
|| jj| jj| jjd}z||   W nF t!y } z,t"j#j$rb t%| j&j'| j(d |W Y d }~n
d }~0 0 |j)}|| j|	< n
| j|	 }| j*|	}t+|}| j,|||tj-krd S fddt./ D }t0t1||gS )Nc                    s   g | ]} | qS r'   r'   rn  r   r'   r(   r.     r/   z2CodeGenerator.call_JitFunction.<locals>.<listcomp>c                 S   s   t |S rB   rI   r*   r   r'   r'   r(   r     r/   z0CodeGenerator.call_JitFunction.<locals>.<lambda>c                    s   i | ]}|t  |qS r'   r   r   r   r'   r(   r     r/   z2CodeGenerator.call_JitFunction.<locals>.<dictcomp>c                 S   s
   t | S rB   r  r  r'   r'   r(   r     r/   c                    s   g | ]}t  |qS r'   r  r   r   r'   r(   r.     r/   c                 S   s   g | ]
}|j qS r'   rd   r  r'   r'   r(   r.     r/   c                 S   s6   g | ].}|d u s$t |tttjjfr,tjjn|jqS rB   )rC   r   r  r   rG   rH   r   r%   r  r'   r'   r(   r.     s
   )
r   r   r   r   rQ   r   r   r   r   r   c                    s   g | ]}  |qS r'   r  r2   )call_opr'   r(   r.     r/   )2inspectgetcallargsrW   rU   rR   rC   r   rH   r  r  r   r   rG   r   r   r?   r   r   Zhas_functionr   r   r  r   r   get_capture_scoper   rQ   r|   r   r   r   r   parse	Exceptionr   r  front_end_debuggingr   r   rS   r   r   Zget_functionrn   callr  r   Zget_num_resultsnextrr   )ru   rW   rX   r   r3   rZ   Zargs_cstZ	args_pathZargs_valfn_namer   r   r   r   rt   eZcallee_ret_typesymbolrm   r'   )rX   r  r(   call_JitFunction  sL    
*
zCodeGenerator.call_JitFunctionc              
      s$  t  |j}t|ts8 j|}|d ur8| |S t|dd}|rt|ddrdt	|j g}t|t
rz|| t jj|d|t fdd|jD } fdd	|jD }ttjd
d |D }t|tr|d|j |j}t|trt|||  |||S t|dr2t|js@tj !|rd j"i}t#$|}	d|	j%v rh |d< z.||i ||}
t|
t&rt&|
}
|
W S  t'y } z*t(j)j*r t jj|d |W Y d }~n
d }~0 0 | j+, v rt-t |}||i |}
t.t/|
r t0|
dd S |
S )NZ_must_use_resultF
_is_unusedz#The result of %s is not being used.r  c                 3   s   | ]}  |V  qd S rB   r   r  r   r'   r(   r     r/   z+CodeGenerator.visit_Call.<locals>.<genexpr>c                    s   g | ]}  |qS r'   r   r  r   r'   r(   r.     r/   z,CodeGenerator.visit_Call.<locals>.<listcomp>c                 s   s"   | ]}t |tr|n|gV  qd S rB   )rC   rM   r   r'   r'   r(   r      r/   r   r   rM  r  c                 S   s   | S rB   r'   )r   r'   r'   r(   r      r/   z*CodeGenerator.visit_Call.<locals>.<lambda>)1r   r   r   rC   r    statically_implemented_functionsr   r   r   r  r  rd  r   r   rS   r;   r  r  rX   rM   	itertoolschainfrom_iterableinsertr   r   r   r[   r  r^   rE   r   rG   
is_builtinr   r  	signature
parametersrN   r  r   r  r  r   rk   mapr`   r%   rj   )ru   rV   rW   Zstatic_implementationZmurerror_messagekwsrX   extra_kwargssigr>   r  r'   r   r(   r     sL    




&



(
zCodeGenerator.visit_Callc                 C   s
   t |jS rB   )r   rg   r   r'   r'   r(   visit_Constant"  s    zCodeGenerator.visit_ConstantrV   c              	   C   sB  | j t|j}|d u r0| |d|jjg }|jD ]}| |}t	|st
|}|du rp|dkrp|  S |du r|dkr|  S q:|j rt|dd }|d ur|| j7 }tjdt| j|t|d || q:t|d	kr|| t|d
kr(| }| }	| ||	|}
||
 qt|dks:J |d	 S )Nz9AST boolean operator '{}' is not (currently) implemented.Flogical_andT
logical_orlinenozeLogical operators 'and' and 'or' are deprecated for non-scalar tensors; please use '&' or '|' instead)categoryfilenamer  sourcer   r   r   )_method_name_for_bool_opr   r%   rF  r   rT  rT   rk   r   rF   r   rK   r   r   r  warn_explicitUserWarningr   r   r  rd  rp   r  rP  )ru   rV   rO  Znontrivial_valuesZsubnoderg   Zbvr  rH  rG  resr'   r'   r(   visit_BoolOp%  sF    





zCodeGenerator.visit_BoolOpr	  r
  r  c                 C   sr   |  |j}t|r,|jdkr,| j|dS t|trF|jdkrF|j}t||j}t	|rnt|t
rnt||S |S )NT)r   r   rg   )r   rg   rF   r   r   ZpermuterC   r   r   rE   r   r   )ru   rV   rG  r   r'   r'   r(   r   [  s    
zCodeGenerator.visit_Attributec                 C   s   d|j _tj| | d S r   )rg   r  r   r  r   r   r'   r'   r(   
visit_Exprg  s    zCodeGenerator.visit_Exprc                 C   s   d S rB   r'   r   r'   r'   r(   visit_NoneTypek  s    zCodeGenerator.visit_NoneTypec                 C   s   t |j}t|D ]\}}t|tjr6t|j||< qt|tjr|j	}| 
|j}t|st| |dtt| |dk rdndt| d |j||< qtdt|qd|S )Nz^Cannot evaluate f-string containing non-constexpr conversion values, found conversion of type r   z{}z{!}z:encountered unexpected node of type {} in a JoinedStr noder   )rM   rk   rR   rC   r   Constantr  rg   FormattedValue
conversionr   rI   r   r%   chrrT  AssertionErrorr;   )ru   rV   rk   r3   rg   Zconversion_codeZ	evaluatedr'   r'   r(   visit_JoinedStrn  s"    

*zCodeGenerator.visit_JoinedStrc                    s*  |d u rd S t    t dt t dt | j}| j }|| _t|drt|dr| j	| j
| j|j |j | j }zt |}W nX ty    Y nF ty } z.tjjr t| jj| jt|d W Y d }~n
d }~0 0 |r|| _| j	| |W  d    S 1 s0    Y  d S )Nignorer  
col_offset)r  catch_warningssimplefilterDeprecationWarningPendingDeprecationWarningr   r|   r  r^   r   r   r   r  r  superr   r   r  r   r  r  r   rS   r1   )ru   rV   Z	last_noder  r>   r  	__class__r'   r(   r     s.    

.zCodeGenerator.visitc                 C   s   |  |dt|jd S )Nzunsupported AST node type: {})r   rT  r%   rT   r   r'   r'   r(   r     s    zCodeGenerator.generic_visitr   c              
   C   s   t |j}d|  k rdkr,n n
t |jr4tdt| |jd }t|tsZtd|s|dkrld}nJz| |jd }W n4 t	y } zdt
| d }W Y d }~n
d }~0 0 t| jj|t|d S )	Nr   r   z=`static_assert` requires one or two positional arguments onlyzqAssertion condition could not be determined at compile-time. Make sure that it depends only on `constexpr` valuesr   r   z'<failed to evaluate assertion message: >)rp   rX   r  r   r   r   rC   r   NotImplementedErrorr  r1   r    r   rS   )ru   rV   	arg_countZpassedr   r  r'   r'   r(   execute_static_assert  s"    
"
&z#CodeGenerator.execute_static_assertc                    s   t jd fdd}|S )Nr  c                    sD   dd  fdd|j D D } fdd|jD }t|i |S )Nc                 S   s   i | ]\}}|t |qS r'   )r   )r,   r$   rg   r'   r'   r(   r     s   z>CodeGenerator.static_executor.<locals>.ret.<locals>.<dictcomp>c                 3   s   | ]}  |V  qd S rB   r   r  r   r'   r(   r     r/   z=CodeGenerator.static_executor.<locals>.ret.<locals>.<genexpr>c                    s   g | ]}t  |qS r'   )r   r   r  r   r'   r(   r.     r/   z>CodeGenerator.static_executor.<locals>.ret.<locals>.<listcomp>)r  rX   r   )ru   rV   r  rX   	python_fnr   r(   r>     s
    z*CodeGenerator.static_executor.<locals>.ret)r   r   )r,  r>   r'   r+  r(   static_executor  s    zCodeGenerator.static_executorr   )r   r  )NFNFNr   )rT   r   r   r   r   r   r  rv   rp   rM   r   r  r  rC   r   r^   r   r   r   updater   rG   Zdevice_printZminimummaximumr   r   r   r
   r   r   r  r  r  r  r   r  r   r$  r   r8  r9  r;  r>  r   r   r   rI  rJ  rK  rP  rU  r   AddSubMultDivFloorDivModPowLShiftRShiftBitAndBitOrBitXorrS  r	   operatorrm  r}  r  r   r   r  r  EqNotEqLtLtEGtGtEr  cmpopr  USubUAddNotInvertr  unaryopr  r  r  r  r  r=  r  r  r  r  r  r   r  r  r  r   r  BoolOpr  AndOrr  boolopr   r  r  r  r   r   r   r*  r-  Zexperimental.gluonZttglZstatic_assertZstatic_printr   r  r   r   __classcell__r'   r'   r%  r(   r     s   
  ;
3	5
4'"5$	N /24&
r   c                    s   d gt  j }j D ] \}}	 j|}
t|	||
< qtg |jj}t	 \}}ddl
m} tdd j} fdd|D }j}|ddd	g||}t||   | d
||||||d}|   |j}||_|S )Nr   )
namedtuplec                 S   s   t | dkS rJ   )rp   )rb   r'   r'   r(   r     r/   zast_to_ttir.<locals>.<lambda>c                    s"   i | ]} j |d   j| qS )r   )rU   r5   r2   rW   rS   r'   r(   r     r/   zast_to_ttir.<locals>.<dictcomp>ZSpecializationProxyr5   r   T)
r   r   r   r   r   r   r   r   r   r   )rp   rU   r   r   indexr   r   r5   r   r   collectionsrN  filterr   r  r1   r   r  r   r   )rW   rS   r   r   r   r   r   r   r   rb   rY   r   r   r   rN  Zleavesr5   r   proxyrt   r>   r'   rO  r(   ast_to_ttir  s&    
rT  )N)Gr   rx   r  r"   r  r  r  Zdataclassesr   ri   r   typingr   r   r   r   r   r	   r
   r   r   r   r   r   Z_C.libtritonr   r   r   r   r   Zlanguage.corer   r   r   Zruntime.jitr   r   Zruntimer   _utilsr   r   r   errorsr   r    r!   r)   r?   r   rE   rF   rI   rL   rO   r[   r`   rj   rn   rg   rr   r  r%   r  rs   r  r   r   r   r   rT  r'   r'   r'   r(   <module>   s\   ,
W:         ?