o
    oh                     @   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
mZmZmZmZmZmZmZmZ ddlmZ ddlmZ ddlmZmZmZmZ ddlmZmZ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)m*Z*m+Z+ dd Z,dd Z-dd Z.de
de/fddZ0de
de/fddZ1de
de/fddZ2de
de/fddZ3de
de/fddZ4d d! Z5d"d# Z6d$d% Z7d&ee fd'd(Z8d)eej9 d*ee fd+d,Z:e/e;e<dhZ=G d-d. d.Z>G d/d0 d0e j?Z@G d1d2 d2ZAG d3d4 d4e j?ZBd5d6 ZCdS )7    N)
ModuleType)	AnyCallableDictOptionalTupleTypeUnionIterableList   )language)ir)	constexprsemantic	str_to_tytensor)_unwrap_if_constexprnv_tma_desc_type
base_value	base_type)get_jit_fn_file_line)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%   r/var/www/html/construction_image-detection-poc/venv/lib/python3.10/site-packages/triton/compiler/code_generator.pycheck_identifier_legality   s   r'   c                 C   s   |   rddtt| j d S |  rdt| j S |  r5tj	j
j}| j|kr,dnd}|t| j S |  r=t| S |  rWt| j}dtt| j}| d| dS |  r]dS td|  )	NT_PiuSVUnsupported type )is_tuplejoinmap	mangle_tytypesis_ptr
element_tyis_intr   dtype
SIGNEDNESSSIGNEDint_signednessstrint_bitwidthis_floatingis_blockscalarshapeis_void	TypeError)tyr:   prefixeltrA   r%   r%   r&   r3      s"   

r3   c                    st   d dd |D }d  fddt D }|dd}|dd}|d	dd
d}|  d| d| }|S )Nr)   c                 S   s   g | ]}t |qS r%   )r3   .0rD   r%   r%   r&   
<listcomp>4       zmangle_fn.<locals>.<listcomp>c                    s"   g | ]}| d t  |  qS )c)reprrH   r+   	constantsr%   r&   rI   5      " ._d_'_sq_[]__)r1   sortedreplace)r"   arg_tysrO   mangled_arg_namesmangled_constantsretr%   rN   r&   	mangle_fn2   s   r^   oreturnc                 C   
   t | tS N)
isinstancer   r_   r%   r%   r&   _is_triton_value>      
re   c                 C   ra   rb   )rc   r   rd   r%   r%   r&   _is_triton_tensorB   rf   rg   c                 C   s   | d u pt | ttjjfS rb   )rc   r   r   corer8   rd   r%   r%   r&   _is_constexprF   s   ri   c                 C   s    t | o| j  p| jjdkS Nr   )rg   r#   r?   numelrd   r%   r%   r&   _is_triton_scalarJ   s    rl   c                 C   s   t | ttfS rb   )rc   listtuplerd   r%   r%   r&   _is_list_likeN   s   ro   c              
   C   sX   |j r(t|D ]"\}}t|s't|s't|j| d|j d|j|  d| qd S 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_argsR   s   r{   c                 C   s   t | tot| tot| dS )N_fields)rc   r#   
issubclassrn   hasattr)valr%   r%   r&   _is_namedtuple\   s   r   c                    sp   t t| r
| j}nt| tjr| jj}n	J dt|   fdd| D }dd |D }t|t||S )NFr/   c                       g | ]} |qS r%   r%   rH   vrw   r%   r&   rI   h   rJ   z*_apply_to_tuple_values.<locals>.<listcomp>c                 S      g | ]}|j qS r%   r#   r   r%   r%   r&   rI   i       )r   r#   r|   rc   r   rn   fields
tuple_type)valuerw   r   valsr4   r%   r   r&   _apply_to_tuple_values`   s   
r   valuesc                 C   s   g }| D ]}| | q|S rb   )_flatten_ir)r   handlesr   r%   r%   r&   flatten_values_to_irm   s   r   r   r4   c                 c   s:    d}|D ]}| | |\}}|V  q|t| ksJ d S )Nr   )_unflatten_irlen)r   r4   cursorrD   r   r%   r%   r&   unflatten_ir_valuest   s   r   c                   @   s$   e Zd Zdd Zdd Zdd ZdS )enter_sub_regionc                 C   
   || _ d S rb   )	generator)selfr   r%   r%   r&   __init__   rf   zenter_sub_region.__init__c                 C   sL   | j j | _| j j | _i | j _| j j | _| j j	 | _
| j| jfS rb   )r   lscopecopyliveins
local_defs	prev_defsbuilderget_insertion_blockinsert_blockget_insertion_pointinsert_pointr   r%   r%   r&   	__enter__   s   zenter_sub_region.__enter__c                 O   s(   | j j| j | j| j _| j| j _d S rb   )r   r   restore_insertion_pointr   r   r   r   r   )r   rx   kwargsr%   r%   r&   __exit__   s   
zenter_sub_region.__exit__N)rt   
__module____qualname__r   r   r   r%   r%   r%   r&   r      s    	r   c                   @   s  e Zd Zdd ZdefddZdefddZdefdd	Zd
ej	defddZ
d
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZdS ) ContainsReturnCheckerc                 C   r   rb   )gscope)r   r   r%   r%   r&   r      rf   zContainsReturnChecker.__init__r`   c                    s   t  fdd|D S )Nc                 3       | ]}  |V  qd S rb   visit)rH   sr   r%   r&   	<genexpr>       z5ContainsReturnChecker._visit_stmts.<locals>.<genexpr>)any)r   bodyr%   r   r&   _visit_stmts   s   z"ContainsReturnChecker._visit_stmtsc                 C   s,   t |tr|js| }t| j|S dS NF)rc   r   rq   parser   r   r   )r   rw   fn_noder%   r%   r&   _visit_function   s   z%ContainsReturnChecker._visit_functionc                 C   sf   d}t |D ])\}}t|tr#|D ]}t|t jr!|p | |}qqt|t jr0|p/| |}q|S r   )astiter_fieldsrc   rm   ASTr   )r   rv   r]   r)   r   itemr%   r%   r&   generic_visit   s   
z#ContainsReturnChecker.generic_visitrv   c                 C   sP   t |jtjr"|jj| jv r | j|jj }t||j}| |S dS | 	|jS r   )
rc   r   r   Nameidr   getattrattrr   r   )r   rv   r   rw   r%   r%   r&   visit_Attribute   s   
z%ContainsReturnChecker.visit_Attributec                 C   s:   t |jtju r
dS |j| jv r| j|j }| |S dS r   )r#   ctxr   Storer   r   r   )r   rv   rw   r%   r%   r&   
visit_Name   s   
z ContainsReturnChecker.visit_Namec                 C      dS )NTr%   r   rv   r%   r%   r&   visit_Return      z"ContainsReturnChecker.visit_Returnc                 C   r   r   r%   r   r%   r%   r&   visit_Assign      z"ContainsReturnChecker.visit_Assignc                 C   r   r   r%   r   r%   r%   r&   visit_AugAssign   r   z%ContainsReturnChecker.visit_AugAssignc                 C      |  |jS rb   r   r   r   r%   r%   r&   visit_Module      z"ContainsReturnChecker.visit_Modulec                 C   r   rb   r   r   r%   r%   r&   visit_FunctionDef   r   z'ContainsReturnChecker.visit_FunctionDefc                 C   s&   |  |j}|jr|p|  |j}|S rb   )r   r   orelse)r   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      z!ContainsReturnChecker.visit_IfExpc                 C   r   rb   )r   funcr   r%   r%   r&   
visit_Call   r   z ContainsReturnChecker.visit_CallN)rt   r   r   r   boolr   r   r   r   	Attributer   r   r   Returnr   Assignr   	AugAssignr   Moduler   FunctionDefr   Ifr   IfExpr   Callr   r%   r%   r%   r&   r      s    r   c                   @   s<   e Zd Zdd ZdejfddZdejfddZdd	 Zd
S )ASTFunctionc                 C   s   || _ || _|| _|| _d S rb   )	ret_types	arg_typesrO   attrs)r   r   r   rO   r   r%   r%   r&   r      s   
zASTFunction.__init__r   c                 C   sH   g }| j D ]}|d u rq||}t|tr|| q|| q|S rb   )r   to_irrc   rm   extendappend)r   r   r   ret_tyir_tyr%   r%   r&   return_types_ir   s   


zASTFunction.return_types_irc                    sF   fdd}t tj|} fdd|D } } ||S )Nc                       |  j vo|d uS rb   rN   pathr)   r   r%   r&   <lambda>  r   z'ASTFunction.serialize.<locals>.<lambda>c                    s   g | ]}t j| qS r%   )r   r   r   rH   r   r   r   r%   r&   rI     s    z)ASTFunction.serialize.<locals>.<listcomp>)rm   r   r   r   get_function_ty)r   r   is_val	val_pathsr   r   r%   r   r&   	serialize   s
   
zASTFunction.serializec              	      s   fdd  j }fdd}ttj |}j D ]\}}|D ]\}}||v r6||||| q$qt|D ]\}	}
tj |
}t	|t
rR||	dd q<t|D ]\}	}
tj |
}t||
t||	| qWj}| D ]\}
}t||
t| qv|S )Nc                    s6   t | tttjfrt fdd| D | S td S )Nc                    r   r%   r%   rH   xmake_templater%   r&   rI     rJ   zBASTFunction.deserialize.<locals>.make_template.<locals>.<listcomp>)rc   rm   rn   r   r   r   )rD   r   r%   r&   r     s   
z.ASTFunction.deserialize.<locals>.make_templatec                    r   rb   rN   r   r   r%   r&   r     r   z)ASTFunction.deserialize.<locals>.<lambda>ztt.nv_tma_descr   )r   rm   r   r   itemsset_arg_attrindexrr   r   rc   r   r   r   r   rx   rO   r   )r   rw   r   r   r   	attr_path
attr_specs	attr_nameattr_valr+   r   rD   rO   r   r%   )r   r   r&   deserialize	  s,   

zASTFunction.deserializeN)	rt   r   r   r   r   r   r   r   r  r%   r%   r%   r&   r      s
    	r   c                       s  e Zd ZU 		ddedee dee fddZd	d
 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dedeeef ddfddZdd Zdd Z dd Z!d d! Z"d"d# Z#d$d% Z$defd&d'Z%d(d) Z&d*d+ Z'd,d- Z(d.d/ Z)d0d1 Z*d2d3 Z+d4d5 Z,d6d7 Z-d8d9 Z.d:d; Z/d<d= Z0d>d? Z1d@dA Z2e3j4dBe3j5dCe3j6dDe3j7dEe3j8dFe3j9dGe3j:dHe3j;dIe3j<dJe3j=dKe3j>dLe3j?dMiZ@eeAe3jB ef edN< dOdP ZCdQdR ZDdSdT ZEdUdV ZFdWdX ZGdYdZ ZHd[d\ ZIe3jJd]e3jKd^e3jLd_e3jMd`e3jNdae3jOdbiZPeeAe3jQ ef edc< ddde ZRe3jSdfe3jTdge3jUdhe3jVdiiZWeeAe3jX ef edj< dkdl ZYdmdn ZZdodp Z[dqdr Z\dsdt Z]dudv Z^dwdx Z_dydz Z`d{d| Zadebeef fd}d~ZcdefddZddefddZedd Zfdd Zgde3jhfddZie3jjde3jkdiZleeAe3jm ef ed< dd Zndd Zodd Zpdd Zq fddZrdd Zsde3jtddfddZudd ZvejjweuejjxeveyeeveeeveiZzee{e|e3jtgef f ed<   Z}S )CodeGeneratorNFr   jit_fnfunction_types	file_namec                 C   sR  || _ t|| _|| _|d | _| j||d || j_|| j_|d u r'i n|| j_|	d u r4| j	 n|	| _
|d u r=i n|| _|| _i | _| D ]/\}}t|tr^||j|| j|< qJt|dd}||v rtt|| |j| j|< qJ|| j|< qJi | _|| _|
rt|d}|| _|
| _d | _|| _g | _d | _i | _|  | _d | _ d| _!d S )Nr   r   r    functionF)"contextr   r   r  
begin_lineset_locoptionscodegen_fns
module_mapcreate_modulemodulefunction_ret_types	prototyper   r   rc   r   getrt   r   r   r  r'   function_name	is_kernelcur_noderq   	scf_stackret_typer   _define_name_lookupdereference_namerw   visiting_arg_default_value)r   r	  r  r   r  r  r  r  r  r  r  r  rq   r  r
  kr   module_namer%   r%   r&   r   )  sD   




zCodeGenerator.__init__c                 C   s   i | ]}|j |qS r%   rt   rH   r)   r%   r%   r&   
<dictcomp>\  rJ   zCodeGenerator.<dictcomp>builtin_namespaceprintminmaxc                 C   s   t | jj||S rb   )r   r  rs   )r   rv   messager%   r%   r&   _unsupportedc  s   zCodeGenerator._unsupportedc                 C   s0   t  }| j||}||u rdS t|rdS dS )NFT)objectr   r  ri   )r   r"   absent_markerr   r%   r%   r&   _is_constexpr_globalf  s   z"CodeGenerator._is_constexpr_globalc                    sJ   dt ffdddt ffddt  dt dtf fdd}|S )	Nr"   c                    s    j | |S rb   )r   r  )r"   absentr   r%   r&   local_lookups  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	dt|t
jt| |  jtjdddkgr@|S ttd	|  d
dd)N__triton_builtin__Fr   r  ztriton.language"TRITON_ALLOW_NON_CONSTEXPR_GLOBALS01z.                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   r8   r   r)  r  osenviron	NameErrortextwrapdedentrY   )r"   r*  r   r   r%   r&   global_lookupw  s(   


z8CodeGenerator._define_name_lookup.<locals>.global_lookupr`   c                    s@    }j jfD ]}|| |}||ur|  S q	t|  d)Nz is not defined)r!  r  r5  )r"   r*  lookup_functionr   r(  r8  r+  r   r%   r&   name_lookup  s   
z6CodeGenerator._define_name_lookup.<locals>.name_lookup)r<   r'  r   )r   r;  r%   r:  r&   r  q  s
   z!CodeGenerator._define_name_lookupr"   r   r`   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)r   r   )r   r"   r   r%   r%   r&   	set_value  s   
zCodeGenerator.set_valuec                 C   s   | j  }| j  }||fS rb   )r   get_locr   )r   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  )r   r?  r>  r%   r%   r&   _set_insertion_point_and_loc  s   z*CodeGenerator._set_insertion_point_and_locc                 C   s8   t |s|g}|D ]}| | t|tjr d S q	d S rb   )ro   r   rc   r   r   )r   stmtsstmtr%   r%   r&   visit_compound_statement  s   
z&CodeGenerator.visit_compound_statementc                 C      t j| | d S rb   r   NodeVisitorr   r   r%   r%   r&   r        zCodeGenerator.visit_Modulec                    s6     |j}|d u sJ t fdd|jD }|S )Nc                       g | ]}  |qS r%   r   )rH   rF   r   r%   r&   rI         z,CodeGenerator.visit_List.<locals>.<listcomp>)r   r   r   rn   elts)r   rv   r   rK  r%   r   r&   
visit_List  s   zCodeGenerator.visit_Listc                    s    |j}g } fdd  |}|d u rtj}nt|tjjs$J || |j}j	
| jd u r;|_nj|krKtdj d| j	 }j	| d S )Nc                    s:   t | tjrt|  S t | tjttfrt| j	S | S rb   )
rc   r   rn   r   r   intfloatr   	to_tensorr   )r   decayr   r%   r&   rQ    s
   
z)CodeGenerator.visit_Return.<locals>.decayzInconsistent return types:  and )r   r   r   voidrc   rh   r   r   r#   r   r]   r  rC   create_blockset_insertion_point_to_end)r   rv   	ret_valuer   r   post_ret_blockr%   rP  r&   r     s    



zCodeGenerator.visit_Returnc                 C   s$   |  |j}t|tjjsJ |jS rb   )r   r   rc   r   rh   rn   r   r   rv   rx   r%   r%   r&   visit_Starred  s   zCodeGenerator.visit_Starredc              	      s,    |j\}} jr |dt|jjd d d D ]G\}}|jj| d  }|j}|j}tj	|t
 d}	|d u rEtj|	g|d}
ntj|	||d}
z jrSJ d _  |
 W d _qd _w  jrid	nd
} j j} j j j|| j _ j j  j } j j}t||D ]
\}} || q j } j|  |j  j  rJ  j d u sȈ j t!j"krt!j" _  j#g  n1t$ j t!j%r j j& j_'n j g j_' j( j j  j# fdd j) jD   j*  |r j+| d S d S )Nz,nested function definition is not supported.r   r   r   targetsr   )targetr   
annotationTFpublicprivatec                    s   g | ]} j |qS r%   )r   create_poisonrG   r   r%   r&   rI          z3CodeGenerator.visit_FunctionDef.<locals>.<listcomp>),r   rx   rw   r&  rr   defaultsr_  rz   r   r   r   r   	AnnAssignr  r  r  r   r   get_or_insert_functionr  r  rq   	push_backadd_entry_blockr  zipr<  r   set_insertion_point_to_startrD  r   has_terminatorr  r   rS  r]   rc   r   r4   r   
reset_typer   finalizerU  )r   rv   ru   kwarg_namesr+   default_valuearg_noder_  r"   	st_target	init_node
visibilityfn_tyentry
arg_valuesarg_name	arg_value	insert_ptr%   r   r&   r     sP   


$
zCodeGenerator.visit_FunctionDefc                 C   s4   g }|j D ]
}|| |g7 }q| |j}||fS rb   )rx   r   kwarg)r   rv   ru   rz   rn  r%   r%   r&   visit_arguments&  s
   
zCodeGenerator.visit_argumentsc                 C   s   t j| | |jS rb   )r   rG  r   rz   r   r%   r%   r&   	visit_arg-  s   zCodeGenerator.visit_argc                 C   sj   |  |j}|  |j}|  |j}|tkr0|| jv r"t| dt|}|| j|< | j| S | |S )Nz4 is already defined. constexpr cannot be reassigned.)r   r_  r^  r   r   r   
ValueErrorr   )r   rv   r_  r^  r   r%   r%   r&   visit_AnnAssign1  s   



zCodeGenerator.visit_AnnAssignc                 C   s   t |tjr|jjjdksJ | ||S t |tjr<|jjjdks$J t|j	D ]\}}| 
| ||j|  q)d S t |tjsDJ | 
| || d S Nr   )rc   r   	Subscriptr   	__class__rt   visit_Subscript_Storer   rr   rK  r<  r   r   r   )r   r^  r   r+   r"   r%   r%   r&   assignTargetA  s   zCodeGenerator.assignTargetc                    s\    fdd   |j}t|tjr|jgn|j}t|dks$J |d | d S )Nc                    sV   t | tjrt|  S tjtjf}t| } | d ur)t| s)t | |s)t| j	} | S rb   )
rc   r   rn   r   r8   r   re   r   rO  r   )r   native_nontensor_types_sanitize_valuer   r%   r&   r  O  s   
z3CodeGenerator.visit_Assign.<locals>._sanitize_valuer   r   )	r   r   rc   r   re  r^  r]  r   r  )r   rv   r   r]  r%   r  r&   r   M  s
   zCodeGenerator.visit_Assignc                 C   sR   |j j}tj|t d}t||j|j}tj|j g|d}| 	| | 
|S )Nr[  r\  )r^  r   r   r   LoadBinOpopr   r   r   r  )r   rv   r"   lhsrhsassignr%   r%   r&   r   _  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   g  s   zCodeGenerator.visit_Namec                 C   rE  rb   rF  r   r%   r%   r&   visit_Storel  rH  zCodeGenerator.visit_Storec                 C   rE  rb   rF  r   r%   r%   r&   
visit_Loado  rH  zCodeGenerator.visit_Loadc                    s    fdd|j D }t|S )Nc                    rI  r%   r   r   r   r%   r&   rI   s  rJ  z-CodeGenerator.visit_Tuple.<locals>.<listcomp>)rK  r   rn   rX  r%   r   r&   visit_Tupler  s   
zCodeGenerator.visit_Tuplec                 C   sT   t |rt|||| jdS t |r#tdd|}t|||| jdS t|||S )N_builderz__(.*)__z__r\1__)rg   r   r   r    sub)r   method_namer  r  reverse_method_namer%   r%   r&   _apply_binary_methodv  s   z"CodeGenerator._apply_binary_methodc                 C   sV   |  |j}|  |j}| jt|j}|d u r$| |d|jj	| 
|||S )Nz8AST binary operator '{}' is not (currently) implemented.)r   leftright_method_name_for_bin_opr  r#   r  r&  formatrt   r  r   rv   r  r  r  r%   r%   r&   visit_BinOp  s   zCodeGenerator.visit_BinOp__add____sub____mul____truediv____floordiv____mod____pow__
__lshift__
__rshift____and____or____xor__r  c                 C   s  | j | | |j | j  }| j }i }|jr9| j | | | _i | _| |j | j }| j  }g }|D ]h}|df|dffD ]4\}	}
||	v r{t	|	| t	|| k}|rg|	| j	|| j	ks{J d| d||  d|
 d|	|  qG||v s||v r|
| ||v r||vr|| ||< ||v r||vr|| ||< q=t| | @ D ]6}||v rq|| }|j	}|| }|j	}t	|t	|k}|r||ksJ d| d| d	| d
|
| q|||||fS )N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   rj  rD  r   r   r   r   r   r   r#   r   rX   keys)r   rv   r   
then_block
else_block	then_defs	else_defsnamesr"   defs
block_name
type_equalthen_valthen_tyelse_valelse_tyr%   r%   r&   visit_then_else_blocks  s^   





z$CodeGenerator.visit_then_else_blocksc                    s  t | }|\}}| j }| j }| j| | j|j|| | ||||\ }}}| j | j| | rCJ | tfdd|D }	| j	|	 | j| | rdJ | t fdd|D }
| j	|
 t
|	t
|
ksJ t|	|
D ]\}}| }|| ksJ | qW d    n1 sw   Y  | j fddtt
|	D }fdd|D }t||}t||D ]
\}}| || qd S )Nc                 3       | ]} | V  qd S rb   r%   rH   r"   r  r%   r&   r         z3CodeGenerator.visit_if_top_level.<locals>.<genexpr>c                 3   r  rb   r%   r  r  r%   r&   r     r  c                    rI  r%   rz   rM   )endif_blockr%   r&   rI     rJ  z4CodeGenerator.visit_if_top_level.<locals>.<listcomp>c                       g | ]} | j qS r%   r   r  r  r%   r&   rI     rJ  )r   r   rT  rU  create_cond_branchhandler  rk  r   create_branchr   ri  get_typeadd_argumentrj  ranger   r<  )r   condrv   srr   ip_blockr  r  r  then_handleselse_handlesthen_helse_hrD   res_handlesr4   
new_valuesr"   	new_valuer%   )r  r  r  r&   visit_if_top_level  s>   




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re| j| |jsm }	n|	
  | j  t|
dkrt fdd|
D }| j| W d    n1 sw   Y  fddtt|D }fd	d|
D }t||}t|
|D ]
\}}| || qd S )
Nc                 3   r  rb   r%   r  r  r%   r&   r     r  z-CodeGenerator.visit_if_scf.<locals>.<genexpr>c                 S      g | ]}|  qS r%   r  rH   hr%   r%   r&   rI     rJ   z.CodeGenerator.visit_if_scf.<locals>.<listcomp>Tr   c                 3   r  rb   r%   r  r  r%   r&   r     r  c                    rI  r%   
get_resultrM   )if_opr%   r&   rI     rJ  c                    r  r%   r   r  r  r%   r&   rI     rJ  )r   r@  r   rT  r   r  r   rA  create_if_opr  merge_block_beforeget_then_blockrU  r   create_yield_opget_else_blockr  r   ri  r<  )r   r  rv   r  r   r)   r?  last_locr  r  r  r  r  r  r4   r  r"   r  r%   )r  r  r  r&   visit_if_scf  s:   



zCodeGenerator.visit_if_scfc              	   C   s   |  |j}t|r6|jtj| jd}t| j |}|r.| j	r&| 
|d| || d S | || d S t|}t|tvrU| 
|dddd tD t|j|rZ|jn|j}| | d S )Nr  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       | ]}|j V  qd S rb   r  r  r%   r%   r&   r   &      z)CodeGenerator.visit_If.<locals>.<genexpr>)r   testrg   tor   int1r   r   r   r  r&  r  r  r   r#   _condition_typesr  r1   rt   r   r   rD  )r   rv   r  contains_returnactive_blockr%   r%   r&   r     s*   zCodeGenerator.visit_Ifc              	   C   s  |  |j}t|r|jtj| jd}t|  |  \}}| j	 }| j
| t|  |j| j}| j }| j	 }| j
| t|  |j| j}| j }| || |j|jkslJ 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 sw   Y  d S t|}t|tvr|  |d!d"dd	 tD t|j#|r|  |jS |  |jS )
Nr  zATernary expression with dynamic condition has inconsistent types rR  Tr   r  r  c                 s   r  rb   r  r  r%   r%   r&   r   Z  r  z,CodeGenerator.visit_IfExp.<locals>.<genexpr>)$r   r  rg   r  r   r  r   r   r@  rT  rj  r   rO  r   r   r   rA  r#   rS  r   r  r  r  r  rU  r  r  rh   r   r  r   r  r&  r  r1   rt   )r   rv   r  r?  r  r  r  r  r  r  ret_type_irr  r%   r%   r&   r   ,  sT   




$#zCodeGenerator.visit_IfExpc                 C      d S rb   r%   r   r%   r%   r&   
visit_Passa  r   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 r:t
||u S t|jd tju rJt
||uS | jt|jd }|d u rf| |d|jd j| |||S )Nr   z1simultaneous multiple comparison is not supportedr   z<AST comparison operator '{}' is not (currently) implemented.)r   comparatorsopsr&  r   r  r   r#   r   Isr   IsNot_method_name_for_comp_opr  r  rt   r  )r   rv   r  r  	lhs_value	rhs_valuer  r%   r%   r&   visit_Compared  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|r,t	||| j
dS zt	|| W S  tyI   | |d| dt|j w )NzAST unary operator 'z!' is not (currently) implemented.r  z)' is not (currently) implemented on type )r   operand_method_name_for_unary_opr  r#   r  r&  rt   rg   r   r   AttributeError)r   rv   r  rw   r%   r%   r&   visit_UnaryOpy  s   zCodeGenerator.visit_UnaryOp__neg____pos____not__
__invert__r  c                 C   s   t |sJ d| dt |sJ d| dt|t|u s(J d| dt|rB|j|jksDJ d| d|j d|j d	d S 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#   rg   )r   r"   loop_vallive_valr%   r%   r&   _verify_loop_carried_variable  s    z+CodeGenerator._verify_loop_carried_variablec                    s  t | &}|\}}|  \}}| j }| j| | j| | |j | j	  | j
}|  g }	g }
|D ]}||v rX|| }|| }| ||| |	| |
| q9t|
}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 s/w   Y  fddtt|D }t||}t|	|D ]\}}|| j|< || j
|< qK|jD ]}J dd S )	Nc                 S   r  r%   r  r  r%   r%   r&   rI     rJ   z-CodeGenerator.visit_While.<locals>.<listcomp>c                 S   r   r%   r   )rH   ar%   r%   r&   rI     r   c                    rI  r%   r  rM   )before_blockr%   r&   rI     rJ  c                    rI  r%   r  rM   )after_blockr%   r&   rI     rJ  c                    rI  r%   r  rM   )while_opr%   r&   rI     rJ  FzNot implemented)#r   r@  r   rT  rj  r  r   rD  r   popr   eraser  r   rA  create_while_opcreate_block_with_parent
get_beforer  r   r   ri  r   r   r  rU  create_condition_opr  	get_afterr   r  r   r   rG  r   )r   rv   r  r   r   r?  r  dummy	loop_defsr  	init_argsr"   r  r  init_handlesinit_tysinit_fe_tys
block_argscondition_argsr   r  body_handles	body_argsyieldsresult_handlesresult_valsnew_defrC  r%   )r  r  r	  r&   visit_While  sz   









C


zCodeGenerator.visit_Whilec                 C   sJ   |j jjdks	J | |j}| |j}t|r!|j|| jdS || S )Nr  r  )	r   r  rt   r   r   slicerg   __getitem__r   )r   rv   r  slicesr%   r%   r&   visit_Subscript_Load  s   z"CodeGenerator.visit_Subscript_Loadc                 C   sJ   |j jjdks	J | |j}| |j}t|tjsJ |	|| d S r  )
r   r  rt   r   r   r   rc   r   rn   __setitem__)r   rv   r   r  r"  r%   r%   r&   r    s
   z#CodeGenerator.visit_Subscript_Storec                 C   s
   |  |S rb   )r#  r   r%   r%   r&   visit_Subscript  rf   zCodeGenerator.visit_Subscriptc                    s    fdd|j D S )Nc                    rI  r%   r   )rH   dimr   r%   r&   rI     rJ  z0CodeGenerator.visit_ExtSlice.<locals>.<listcomp>)dimsr   r%   r   r&   visit_ExtSlice  s   zCodeGenerator.visit_ExtSlicec           +         s0   |jj}fdd|jjD }tfdd|jjD }|tjkrZ||i |}t|j	j
|jj
|jj
}|D ]}t|j|jj< |j |jD ]	}tj| qMq9d S d }	d }
d}d}|tju r||i |}|j	}|j}|j}|j}	|j}
|j}|j}n@|tu rt|dkr|d n td}t|dkr|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
}||}}t |j!}t |j!}t |j!}|j"# r|j"# r|j"# st$d|j" d|j" d|j" dt%|j"|j"}t%||j"}|&j!}|j'tj(j"j)j*k}|j+}|j+}|j+}j!,|||}j!,|||}j!,|||}j!-|}.|jjtj(/|| t0v}|\}}1 \}}j!2 }j!3| j45| |j j46  |7  g }g }g }j8D ]'}||v rňj8| }|| } 9|||  |5| |5|  |5| q:|| t;|}!dd |D }"j!<||||! t=|	d ur >dj!?|	 t=|
d ur >dj!?|
 |r >dj!@  |r >dj!@  j45|  Adj!3 |B _i _8fddtt|!D }#tC|#|"}$tD||$D ]\}}%.||% qK|j j46  g }j8D ]}||v rj8| }&tE|&trt |&j!}&|5|& qgt|dkrt;|}'j!F|' G }(|(H dksJ dj!3  I }|rj!J||}j!K||}j|jj j+L| .|jjtj(/|| W d    n	1 sw   Y   fddtt|!D })tC|)|"}*tD||*D ]\}}%.||% q|jD ]}J dd S )Nc                    rI  r%   r   rH   rz   r   r%   r&   rI     rJ  z+CodeGenerator.visit_For.<locals>.<listcomp>c                 3   r   rb   r   rH   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  r  c                 S   r   r%   r   r   r%   r%   r&   rI   W  r   ztt.num_stagesztt.loop_unroll_factorztt.disallow_acc_multi_bufferz
tt.flattenc                    s   g | ]	}  |d  qS )r   r  rM   )for_op_bodyr%   r&   rI   h      z7We use SCF, so the loop body should only have one blockc                    rI  r%   r  rM   )for_opr%   r&   rI     rJ  z)Don't know what to do with else after for)Mr   iterr   rx   dictkeywordsr   static_ranger  startr   endstepr   r   r^  r   rD  r   r   r   rG  r   
num_stagesloop_unroll_factordisallow_acc_multi_bufferflattenr   NumRuntimeErrorri   r   rO  r   r8   r7   rC   integer_promote_implr   r;   rh   r9   r:   r  create_int_castrb  r<  r   r   r@  rT  rj  r  r   r
  r  r   r  rA  r   create_for_opr   set_attrget_int32_attrget_unit_attrget_bodyr   r   ri  rc   r  
get_parentsizeget_induction_var
create_sub
create_addreplace_all_uses_with)+r   rv   IteratorClass	iter_argsiter_kwargsiteratorr2  r+   rC  r6  r7  r8  r9  lbubr5  negative_stepiv_type
iv_ir_typeiv_is_signedivr  r   r   r?  r  blockr  r  r  r"   r  r  r  r  block_handlesr  r   localyield_handlesfor_op_regionr  result_valuesr%   )r.  r,  r   r&   	visit_For  s   


$&&
$"















O

zCodeGenerator.visit_Forc                 C   s2   |  |j}|  |j}|  |j}t|||S rb   )r   lowerupperr5  r   r   )r   rv   r[  r\  r5  r%   r%   r&   visit_Slice  s   zCodeGenerator.visit_Slicec                 C   r   rb   )r   r   r   r%   r%   r&   visit_Index  r   zCodeGenerator.visit_Indexc                 C   s   |j | |jfS rb   )rz   r   r   r   r%   r%   r&   visit_keyword  rH  zCodeGenerator.visit_keywordc                 C   s:   |  |j}|jd ur|  |jnd}tjj||| jdS )Nr  r  )r   r  msgr   rh   device_assertr   )r   rv   r  r`  r%   r%   r&   visit_Assert  s   zCodeGenerator.visit_Assertrw   c                    s  t j|jg R i |  fdd|jD  t D ]\}}t|tjtt	t
tfr2tj| |< qt dd } fdd|D }t dd } fdd|D }t|jd	d |D |}	| j|	s|j}
t|\}}d
d  D }tg ||t }t| j||
| j||	| j|j||| jj| jj| jjd}z	||   W n t!y } z
t"| j#j$| j%d |d }~ww |j&}|| j|	< n| j|	 }| j'|	}dd |D }| j(|||tj)krd S fddt*+ D }t,t-||gS )Nc                    s   g | ]} | qS r%   r%   r  rx   r%   r&   rI     rJ   z2CodeGenerator.call_JitFunction.<locals>.<listcomp>c                 S   s   t |S rb   ri   r)   r   r%   r%   r&   r     s    z0CodeGenerator.call_JitFunction.<locals>.<lambda>c                    s   i | ]}|t  |qS r%   r   r   rc  r%   r&   r     rc  z2CodeGenerator.call_JitFunction.<locals>.<dictcomp>c                 S   s
   t | S rb   rd  re  r%   r%   r&   r     s   
 c                    s   g | ]}t  |qS r%   rf  r   rc  r%   r&   rI     rJ  c                 S   r   r%   r   r)  r%   r%   r&   rI     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   rM  r   rh   r8   r   r#   r)  r%   r%   r&   rI     s    )
r  r  r  r  rq   r  r
  r  r  r  c                 S   r   r%   )r  r)  r%   r%   r&   rI     r   c                    rI  r%   r  rM   )call_opr%   r&   rI     rJ  ).inspectgetcallargsrw   ru   rr   rc   r   r8   rN  rM  r   r   rh   r   r   r^   rt   r  has_function__globals__r   r   r0  r  r	  r  rq   r   r  r  r  r   r   	Exceptionr   r  rs   r  r  get_functioncallrS  r  get_num_resultsnextr   )r   rw   rx   r   r+   rz   args_cst	args_pathargs_valfn_namer   r  r
  r   r  r   ecallee_ret_typesymbolr   r%   )rx   rg  r&   call_JitFunction  sP   

zCodeGenerator.call_JitFunctionc           
   
      sz  t  |j} j|}|d ur| |S t fdd|jD } fdd|jD }tt	j
dd |D }t|trLt|||  |||S t|drVt|js\tj|rd ji}t|}d|jv ro |d< z||i ||}t|trt|}|W S  ty }	 z	t jj|d |	d }	~	ww | j  v rt!t |}||i |}t"t#|rt$|d	d
 S |S )Nc                 3   r   rb   r   r*  r   r%   r&   r     r   z+CodeGenerator.visit_Call.<locals>.<genexpr>c                    rI  r%   r   r)  r   r%   r&   rI     rJ  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     s   " __self__r  
_generatorc                 S   s   | S rb   r%   )r   r%   r%   r&   r     s    z*CodeGenerator.visit_Call.<locals>.<lambda>)%r   r   r    statically_implemented_functionsr  r0  r1  rx   rm   	itertoolschainfrom_iterablerc   r   r{   rx  r~   re   ry  r   rh   
is_builtinr   rh  	signature
parametersrn   rl  r   r  rs   r!  r   r2   r   r#   r   )
r   rv   rw   static_implementationkwsrx   extra_kwargssigr]   ru  r%   r   r&   r     s8   

 




	
zCodeGenerator.visit_Callc                 C   s
   t |jS rb   )r   r   r   r%   r%   r&   visit_Constant  rf   zCodeGenerator.visit_Constantrv   c                 C   sx   t |jdkr| |d| |jd }| |jd }| jt|j}|d u r5| |d|jj	| 
|||S )Nr   z^chained boolean operators (A or B or C) are not supported; use parentheses to split the chain.r   r   z9AST boolean operator '{}' is not (currently) implemented.)r   r   r&  r   _method_name_for_bool_opr  r#   r  r  rt   r  r  r%   r%   r&   visit_BoolOp  s   zCodeGenerator.visit_BoolOplogical_and
logical_orr  c                 C   s<   |  |j}t|r|jdkrtj|d| jdS t||jS )Nr(   )r   r   )r   )r   r   rg   r   r   permuter   r   )r   rv   r  r%   r%   r&   r     s   zCodeGenerator.visit_Attributec                 C   rE  rb   rF  r   r%   r%   r&   
visit_Expr  rH  zCodeGenerator.visit_Exprc                 C   r  rb   r%   r   r%   r%   r&   visit_NoneType  r   zCodeGenerator.visit_NoneTypec                 C   s   t |j}t|D ]N\}}t|tjrt|j||< q	t|tjrO|j	}| 
|j}t|s:| |dtt| |dk r@dndt| d |j||< q	tdt|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   r   rr   rc   r   Constantr<   r   FormattedValue
conversionr   ri   r&  r#   chrr  AssertionErrorr1   )r   rv   r   r+   r   conversion_code	evaluatedr%   r%   r&   visit_JoinedStr  s"   

*
zCodeGenerator.visit_JoinedStrc                    s
  |d u rd S t  q t dt t dt | j}| j }|| _t|dr?t|dr?| j	| j
| j|j |j | j }zt |}W n tyO     tyf } zt| jj| jt|d d }~ww |rr|| _| j	| |W  d    S 1 s~w   Y  d S )Nignorelineno
col_offset)warningscatch_warningssimplefilterDeprecationWarningPendingDeprecationWarningr  r   r=  r~   r  r  r
  r  r  superr   r   rl  r  rs   rL   )r   rv   	last_noder  r]   ru  r  r%   r&   r   #  s0   


$zCodeGenerator.visitc                 C   s   |  |dt|j)Nzunsupported AST node type: {})r&  r  r#   rt   r   r%   r%   r&   r   @  r   zCodeGenerator.generic_visitc              
   C   s   t |j}d|  k rdkrn tdt |jrtdt| |jd }t|ts0td|sh|dkr9d}n%z
| |jd }W n t	y] } zdt
| d }W Y d }~nd }~ww 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: >)r   rx   r1  rC   r   r   rc   r   NotImplementedErrorrl  rL   r   r  rs   )r   rv   	arg_countpassedr%  ru  r%   r%   r&   execute_static_assertC  s*   


z#CodeGenerator.execute_static_assertc                    s   dt jf fdd}|S )Nrv   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   )rH   r"   r   r%   r%   r&   r   \  s    z>CodeGenerator.static_executor.<locals>.ret.<locals>.<dictcomp>c                 3   r   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&   rI   `  r-  z>CodeGenerator.static_executor.<locals>.ret.<locals>.<listcomp>)r1  rx   r   )r   rv   r  rx   	python_fnr   r&   r]   [  s
   z*CodeGenerator.static_executor.<locals>.ret)r   r   )r  r]   r%   r  r&   static_executorY  s   zCodeGenerator.static_executorr{  )NFNFNr   )~rt   r   r   r   r   r   r<   r   r   rm   r  rN  rM  rc   r   r!  r   __annotations__updater   rh   device_printminimummaximumr&  r)  r  r	   r   r   r<  r@  rA  rD  r   rL  r   rY  r   r{  r|  r~  r  r   r   r   r  r  r  r  r  r   AddSubMultDivFloorDivModPowLShiftRShiftBitAndBitOrBitXorr  r   operatorr  r  r  r   r   r  r  EqNotEqLtLtEGtGtEr  cmpopr  USubUAddNotInvertr  unaryopr  r  r#  r  r%  r(  rZ  r]  r^  r   r_  rb  rx  r   r  BoolOpr  AndOrr  boolopr   r  r  r  r   r   r   r  r  static_assertstatic_printr"  r{  r'  r   __classcell__r%   r%   r  r&   r  '  s   
 
,3
0	5	
4'5$	N .%&(r  c                    s   t ttj }tg |jj}t \}}	ddl	m
}
 tdd j} fdd|D }j}|
ddd	g||}t|| j  | d
||	|||d}|   |j}||_|S )Nr   )
namedtuplec                 S   s   t | dkS rj   )r   )r   r%   r%   r&   r   s  s    zast_to_ttir.<locals>.<lambda>c                    s"   i | ]} j |d   j| qS )r   )ru   rO   rM   rw   rs   r%   r&   r   t  rP   zast_to_ttir.<locals>.<dictcomp>SpecializationProxyrO   r  T)	r   r  r  r  r  r
  r  r  r  )rm   r2   r   r  r   r   rO   r   r   collectionsr  filterr  rk  r   rL   r   r   r  r	  )rw   rs   r	  r  r  r  r   r  r  r
  r  leavesrO   r  proxyr   r]   r%   r  r&   ast_to_ttirm  s    r  )Dr   rh  r    r  r3  r6  r|  r4   r   typingr   r   r   r   r   r   r	   r
   r   r  r   _C.libtritonr   r   r   r   r   language.corer   r   r   r   runtime.jitr   runtimer   _utilsr   r   r   errorsr   r   r   r'   r3   r^   r   re   rg   ri   rl   ro   r{   r   r   r   r   r   rM  r#   r  r   rG  r   r   r  r  r%   r%   r%   r&   <module>   sX    ,
X;        N