U
    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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 ddlmZ ddlmZ d	d
lmZmZmZ d dlmZ dd Z dd Z!ee"dddZ#ee"dddZ$ee"dddZ%ee"dddZ&edddZ'dd Z(dd Z)e"e*e+dhZ,G d d! d!Z-G d"d# d#e j.Z/G d$d% d%e j.Z0d&d' Z1d(d) Z2dS )*    N)AnyCallableDictOptionalTupleTypeUnion   )language)ir)	constexprtensor	str_to_ty)_normalize_ty)JITFunction   )CompilationErrorCompileTimeAssertionFailureUnsupportedLanguageConstruct)
ModuleTypec                 C   s   |   rdt| j S |  rHtjjj}| j|kr6dnd}|t	| j
 S |  rXt	| S |  rt| j}dtt	| j}| d| dS |  rdS dstdd S )	NPiu_SVFzUnsupported type)Zis_ptr	mangle_tyZ
element_tyis_intr
   dtype
SIGNEDNESSSIGNEDint_signednessstrZint_bitwidthZis_floatingis_blockZscalarjoinmapshapeZis_voidAssertionError)tyr    prefixeltr&    r+   P/var/www/html/venv/lib/python3.8/site-packages/triton/compiler/code_generator.pyr      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 )Nr   c                 S   s   g | ]}t |qS r+   )r   .0r(   r+   r+   r,   
<listcomp>'   s     zmangle_fn.<locals>.<listcomp>c                    s"   g | ]}| d t  |  qS )c)reprr.   r   	constantsr+   r,   r/   (   s     .Z_d_'Z_sq_[]__)r$   sortedreplace)nameZarg_tysr4   Zmangled_arg_namesZmangled_constantsretr+   r3   r,   	mangle_fn%   s    r>   )oreturnc                 C   s
   t | tS N)
isinstancer   r?   r+   r+   r,   _is_triton_tensor1   s    rD   c                 C   s
   t | tS rA   )rB   r   rC   r+   r+   r,   _is_constexpr5   s    rE   c                 C   s    t | o| j  p| jjdkS )Nr   )rD   typer#   ZnumelrC   r+   r+   r,   _is_triton_scalar9   s    rG   c                 C   s   t | ttfS rA   )rB   listtuplerC   r+   r+   r,   _is_list_like=   s    rJ   rC   c                 C   s   t | tr| jS | S rA   )rB   r   valuerC   r+   r+   r,   _unwrap_if_constexprA   s    rL   c              
   C   sT   |j rPt|D ]@\}}t|st|st|j| d|j d|j|  d| qd S )Nz	Function z= is marked noinline, but was called with non-scalar argument :)noinline	enumeraterE   rG   r   src__name__	arg_names)nodefnargsidxargr+   r+   r,   _check_fn_argsE   s     rX   c                 C   sd   | }t |ts|j}q|jjj}t|j\}}t|D ]"\}}| 	dr8||7 } q\q8||fS )Nzdef )
rB   r   rT   __code__co_filenameinspectgetsourcelinesrO   strip
startswith)rT   Zbase_fn	file_namelines
begin_linerV   liner+   r+   r,   _get_fn_file_lineO   s    

rc   c                   @   s$   e Zd Zdd Zdd Zdd ZdS )enter_sub_regionc                 C   s
   || _ d S rA   )	generator)selfre   r+   r+   r,   __init__f   s    zenter_sub_region.__init__c                 C   sL   | j j | _| j j | _i | j _| j j | _| j j	 | _
| j| jfS rA   )re   lscopecopyliveins
local_defs	prev_defsbuilderget_insertion_blockinsert_blockget_insertion_pointinsert_pointrf   r+   r+   r,   	__enter__i   s    zenter_sub_region.__enter__c                 O   s(   | j j| j | j| j _| j| j _d S rA   )re   rm   restore_insertion_pointrq   rj   rh   rl   rk   )rf   rU   kwargsr+   r+   r,   __exit__r   s    
zenter_sub_region.__exit__N)rQ   
__module____qualname__rg   rs   rv   r+   r+   r+   r,   rd   d   s   	rd   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 rA   )gscope)rf   rz   r+   r+   r,   rg   {   s    zContainsReturnChecker.__init__r@   c                 C   s   |D ]}|  |r dS qdS )NTFvisit)rf   bodysr+   r+   r,   _visit_stmts~   s    
z"ContainsReturnChecker._visit_stmtsc                 C   s,   t |tr(|js(| }t| j|S dS NF)rB   r   rN   parsery   rz   r}   )rf   rT   Zfn_noder+   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_fieldsrB   rH   ASTr}   )rf   rS   r=   r   rK   itemr+   r+   r,   generic_visit   s    
z#ContainsReturnChecker.generic_visitrS   r@   c                 C   sP   t |jtjrD|jj| jkr@| j|jj }t||j}| |S dS | 	|jS r   )
rB   rK   r   Nameidrz   getattrattrr   r}   )rf   rS   rK   rT   r+   r+   r,   visit_Attribute   s    
z%ContainsReturnChecker.visit_Attributec                 C   s:   t |jtjkrdS |j| jkr6| j|j }| |S dS r   )rF   ctxr   Storer   rz   r   )rf   rS   rT   r+   r+   r,   
visit_Name   s    
z ContainsReturnChecker.visit_Namec                 C   s   dS )NTr+   rf   rS   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 rA   r   r~   r   r+   r+   r,   visit_Module   s    z"ContainsReturnChecker.visit_Modulec                 C   s   |  |jS rA   r   r   r+   r+   r,   visit_FunctionDef   s    z'ContainsReturnChecker.visit_FunctionDefc                 C   s&   |  |j}|jr"|p |  |j}|S rA   )r   r~   orelse)rf   rS   r=   r+   r+   r,   visit_If   s    zContainsReturnChecker.visit_Ifc                 C   s   |  |jp|  |jS rA   )r}   r~   r   r   r+   r+   r,   visit_IfExp   s    z!ContainsReturnChecker.visit_IfExpc                 C   s   |  |jS rA   )r}   funcr   r+   r+   r,   
visit_Call   s    z ContainsReturnChecker.visit_CallN)rQ   rw   rx   rg   boolr   r   r   r   	Attributer   r   r   Returnr   Assignr   	AugAssignr   Moduler   FunctionDefr   Ifr   IfExpr   Callr   r+   r+   r+   r,   ry   y   s   ry   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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$d"d# Z%d$d% Z&d&d' Z'd(d) Z(d*d+ Z)d,d- Z*d.d/ Z+d0d1 Z,d2d3 Z-d4d5 Z.d6d7 Z/d8d9 Z0e1j2d:e1j3d;e1j4d<e1j5d=e1j6d>e1j7d?e1j8d@e1j9dAe1j:dBe1j;dCe1j<dDe1j=dEiZ>ee?e1j@ ef edF< dGdH ZAdIdJ ZBdKdL ZCdMdN ZDdOdP ZEdQdR ZFdSdT ZGe1jHdUe1jIdVe1jJdWe1jKdXe1jLdYe1jMdZiZNee?e1jO ef ed[< d\d] ZPe1jQd^e1jRd_e1jSd`e1jTdaiZUee?e1jV ef edb< dcdd ZWdedf ZXdgdh ZYdidj ZZdkdl Z[dmdn Z\e]eef dodpdqZ^edodrdsZ_edtdudvZ`dwdx Zadydz Zbe1jcd{d|d}Zde1jed~e1jfdiZgee?e1jh ef ed< eijjdk rdd Zkdd Zldd Zmdd Zndd Zodd Zpdd Zq fddZrdd Zse1jtddddZudd ZvejjweuejjxeveyeeveeeveiZzee{e|e1jtgef f ed<   Z}S )CodeGeneratorNFr   )jit_fnfunction_typesr_   c                 C   s   || _ t|| _|| _|d | _| j||d || j_|	| j_|d krT| j n|| _	|d krfi n|| _
|| _|| _t | _|| _|| _|| _|| _|| _d | _|
d kr|jn|
| _|| _g | _d | _i | _|  | _d | _d| _d S )Nr   r   F)contextr   rm   r_   ra   set_locoptionscodegen_fnscreate_modulemodulefunction_ret_types	prototyperz   dictrh   
attributesr4   r   function_name	is_kernelcur_nodedebugrN   	scf_stackret_typerk   _define_name_lookupdereference_namerT   visiting_arg_default_value)rf   r   r   rz   r   r4   r   r   r   r   r   r   r   r   rN   r_   ra   r+   r+   r,   rg      s4    

zCodeGenerator.__init__c                 C   s   i | ]}|j |qS r+   rQ   r.   r   r+   r+   r,   
<dictcomp>   s      zCodeGenerator.<dictcomp>builtin_namespaceprintminmaxc                 C   s   t | jj||S rA   )r   r   rP   )rf   rS   messager+   r+   r,   _unsupported  s    zCodeGenerator._unsupportedc                 C   sT   t  }| j||}||kr dS t|r,dS | jdi | }rPt|dkS dS )NFT__annotations__r   )objectrz   getrE   r   )rf   r<   absent_markervalar+   r+   r,   _is_constexpr_global  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 rA   )rh   r   )r<   absentrr   r+   r,   local_lookup  s    z7CodeGenerator._define_name_lookup.<locals>.local_lookupc                    s    j | |}||ks|  jkst|tkst|tst|ddst|dddst|t	j
s | s jstjdddkr|S ttd	|  d
ddd S )NZ__triton_builtin__Frw    ztriton.languageZ"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 annotated as constexpr (`x: triton.language.constexpr = 42`
                or `x = triton.language.constexpr(42)`).  Alternatively, set the
                envvar TRITON_ALLOW_NON_CONSTEXPR_GLOBALS=1, but we do not
                promise to support this forever.
 )rz   r   r   rF   r   rB   r   r   r^   r
   r   r   r   osenviron	NameErrortextwrapdedentr;   )r<   r   r   rr   r+   r,   global_lookup  s:    



 z8CodeGenerator._define_name_lookup.<locals>.global_lookup)r<   r@   c                    sD    }j jfD ]}|| |}||k	r|  S qt|  dd S )Nz is not defined)r   r   r   )r<   r   Zlookup_functionrK   r   r   r   rf   r+   r,   name_lookup8  s    

z6CodeGenerator._define_name_lookup.<locals>.name_lookup)r"   r   r   )rf   r   r+   r   r,   r     s
    z!CodeGenerator._define_name_lookup)r<   rK   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)rh   rk   )rf   r<   rK   r+   r+   r,   	set_valueB  s    
zCodeGenerator.set_valuec                 C   s   | j  }| j  }||fS rA   )rm   get_locrp   )rf   locipr+   r+   r,   _get_insertion_point_and_locK  s    

z*CodeGenerator._get_insertion_point_and_locc                 C   s   | j | | j | d S rA   )rm   rt   r   )rf   r   r   r+   r+   r,   _set_insertion_point_and_locS  s    z*CodeGenerator._set_insertion_point_and_locc                 C   s6   t |s|g}|D ]}| | t|tjr q2qd S rA   )rJ   r}   rB   r   r   )rf   Zstmtsstmtr+   r+   r,   visit_compound_statementZ  s    
z&CodeGenerator.visit_compound_statementc                 C   s   t j| | d S rA   r   NodeVisitorr   r   r+   r+   r,   r   f  s    zCodeGenerator.visit_Modulec                    s0     |j}|d kst fdd|jD }|S )Nc                    s   g | ]}  |qS r+   r|   )r.   r*   rr   r+   r,   r/   l  s     z,CodeGenerator.visit_List.<locals>.<listcomp>)r}   r   r'   elts)rf   rS   r   r   r+   rr   r,   
visit_Listi  s    zCodeGenerator.visit_Listc                    s     |j}|d kr( jg  tj}npt|trr fdd|D }dd |D } jdd |D  t|}n&tj	| j} j|j
g |j} jd kr| _n  j|krtd j d| d S )Nc                    s   g | ]}t j| jqS r+   )r
   core
_to_tensorrm   r.   vrr   r+   r,   r/   z  s     z.CodeGenerator.visit_Return.<locals>.<listcomp>c                 S   s   g | ]
}|j qS r+   rF   r   r+   r+   r,   r/   {  s     c                 S   s   g | ]
}|j qS r+   handler   r+   r+   r,   r/   |  s     zInconsistent return types:  and )r}   rK   rm   r=   r
   voidrB   rI   r   r   r   rF   r   	TypeError)rf   rS   Z	ret_valueZret_tyZ
ret_values	ret_typesr=   r+   rr   r,   r   p  s     



zCodeGenerator.visit_Returnc              	   C   s  |  |j\}}| jr"| |dt|jjD ]\}}|jj| d  }|j}|j}tj	|t
 d}	|d krtj|	g|d}
ntj|	||d}
z| jrtd| _|  |
 W 5 d| _X q.| jrdnd	}| j| j| j| j| j|| j| _| j| j | j }g }d
}t|D ]\}}|| jkr\| j| }t|sLt| j| }|| qnZ|| jkr| j| D ]\}}| j||| qr|t| j|| jj |  |d7 }q| j! }t"||D ]\}}| #|| q| j$| | %|j& | j'd ks| j't(j)kr.t(j)| _'| j*g  nVt+| j't,rbt-| j'| j_.| j/| j| j n"| j'g| j_.| j/| j| j |r| j0| | j1  d S )Nz,nested function definition is not supported.r   r   r   targetsrK   )targetrK   
annotationFTpublicprivater   )2r}   rU   rT   r   rO   defaultsr   rW   r   r   r   r   	AnnAssignr   r'   r   rm   Zget_or_insert_functionr   r   r   to_irrN   Z	push_backZadd_entry_blockr4   rE   r   appendr   Zset_arg_attrr   Zparam_typesrn   zipr   set_insertion_point_to_startr   r~   r   r
   r   r=   rB   rI   rH   r   Z
reset_typeset_insertion_point_to_endfinalize)rf   rS   rR   kwarg_namesr   default_valueZarg_noder   r<   Z	st_targetZ	init_nodeZ
visibilityentryZ
arg_valuesrV   arg_nameZcstrK   Z	insert_pt	arg_valuer+   r+   r,   r     sl    

  



 
zCodeGenerator.visit_FunctionDefc                 C   s4   g }|j D ]}|| |g7 }q
| |j}||fS rA   )rU   r}   kwarg)rf   rS   rR   rW   r  r+   r+   r,   visit_arguments  s
    
zCodeGenerator.visit_argumentsc                 C   s   t j| | |jS rA   )r   r   r   rW   r   r+   r+   r,   	visit_arg  s    zCodeGenerator.visit_argc                 C   sr   |  |j}|  |j}|  |j}|tkrh|| jkrDt| dt|sTt|}|| j|< | j| S | |S )Nz4 is already defined. constexpr cannot be reassigned.)	r}   r   r   rK   r   rh   
ValueErrorrE   r   )rf   rS   r   r   rK   r+   r+   r,   visit_AnnAssign  s    


zCodeGenerator.visit_AnnAssignc           	      C   s   g }|j D ]}|| |g7 }q
t|dkr8| |d|d }| |j}t|sZ|g}t|sh|g}tjf}t||D ]F\}}t	|}|d k	rt
|st||stj|| j}| || qzd S )Nr   z2simultaneous multiple assignment is not supported.r   )r   r}   lenr   rK   rJ   r
   r   r   rL   rD   rB   r   r   rm   r   )	rf   rS   Z_namesr   namesvaluesZnative_nontensor_typesr<   rK   r+   r+   r,   r     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oprK   r   r}   r   )rf   rS   r<   lhsrhsZassignr+   r+   r,   r     s    
zCodeGenerator.visit_AugAssignc                 C   s"   t |jtjkr|jS | |jS rA   )rF   r   r   r   r   r   r   r+   r+   r,   r     s    zCodeGenerator.visit_Namec                 C   s   t j| | d S rA   r   r   r+   r+   r,   visit_Store  s    zCodeGenerator.visit_Storec                 C   s   t j| | d S rA   r   r   r+   r+   r,   
visit_Load  s    zCodeGenerator.visit_Loadc                    s    fdd|j D }t|S )Nc                    s   g | ]}  |qS r+   r|   )r.   xrr   r+   r,   r/     s     z-CodeGenerator.visit_Tuple.<locals>.<listcomp>)r   rI   )rf   rS   rU   r+   rr   r,   visit_Tuple  s    zCodeGenerator.visit_Tuplec                 C   sT   t |rt|||| jdS t |rFtdd|}t|||| jdS t|||S )NZ_builderz__(.*)__z__r\1__)rD   r   rm   resub)rf   method_namer  r  Zreverse_method_namer+   r+   r,   _apply_binary_method  s    z"CodeGenerator._apply_binary_methodc                 C   sV   |  |j}|  |j}| jt|j}|d krH| |d|jj	| 
|||S )Nz8AST binary operator '{}' is not (currently) implemented.)r}   leftright_method_name_for_bin_opr   rF   r  r   formatrQ   r  rf   rS   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rr| j | | | _i | _| |j | j }| j  }g }g }g }	|D ]}
|df|dffD ]P\}}|
|kr||
 j	||
 j	kst
d|
 d||
 j	 d| d||
 j	 q|
|ks|
|krV||
 ||
|kr||
 j	n||
 j	 |	|
|krF||
 j n||
 j  |
|krv|
|krv||
 ||
< |
|kr|
|kr||
 ||
< q| | @ D ]v}
|
|krq||
 j	}||
 j	}||kst
d|
 d| d	| d
||
 || |	||
 j  q|||||||	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 ())rm   r   r   r~   rn   rk   ri   r   rh   rF   r'   r   r   get_typekeys)rf   rS   rj   
then_block
else_block	then_defs	else_defsr  r   ir_ret_typesr<   ZdefsZ
block_nameZthen_tyZelse_tyr+   r+   r,   visit_then_else_blocks7  sZ    





&
$




z$CodeGenerator.visit_then_else_blocksc              	      sj  d}t | }|\}}| j }| j }| j }	| j| | j|j|| | ||||\ }}}
}}| j| | r| rd}|	  |	 s|r| j
|	fdd|
D  | j| |	 s|r| j
|	 fdd|
D  |r|D ]}|	| q W 5 Q R X |rf| j|	 t|
D ].\}}tj|	||| }| || q6d S )NTFc                    s   g | ]} | j qS r+   r   r.   nr6  r+   r,   r/     s     z4CodeGenerator.visit_if_top_level.<locals>.<listcomp>c                    s   g | ]} | j qS r+   r   r:  r7  r+   r,   r/     s     )rd   rm   create_blockr  Zcreate_cond_branchr   r9  Z
has_returneraseZhas_terminatorZcreate_branchadd_argumentr   rO   r
   r   r   rW   r   )rf   condrS   Zhas_endif_blocksrrj   Zip_blockr4  r5  Zendif_blockr  r   r8  r(   r   r<   
new_tensorr+   )r7  r6  r,   visit_if_top_leveln  s6    


z CodeGenerator.visit_if_top_levelc              	      sh  t }|\}} \}}j }|jr:j nd }	||||	\ }}	}
}}|| jfdd|D |jd}|	|
  j|
  t|
dkr̈jfdd|
D  |js| }	n|		|  j|  t|
dkr"j fdd|
D  W 5 Q R X t|
D ].\}}tj|||| }|| q4d S )Nc                    s   g | ]}|  jqS r+   r   rm   r-   rr   r+   r,   r/     s     z.CodeGenerator.visit_if_scf.<locals>.<listcomp>Tr   c                    s   g | ]} | j qS r+   r   r:  r<  r+   r,   r/     s     c                    s   g | ]} | j qS r+   r   r:  r=  r+   r,   r/     s     )rd   r   rm   r>  r   r9  r   create_if_opr   merge_block_beforeget_then_blockr  r  create_yield_opget_else_blockrO   r
   r   r   
get_resultr   )rf   rA  rS   rB  rj   r   r   last_locr4  r5  r  r   if_opr   r<   rC  r+   )r7  rf   r6  r,   visit_if_scf  s,    
 
$zCodeGenerator.visit_if_scfc              	   C   s   |  |j}t|rt|jtj| jd}t| j |}| j	rN|rN| 
|dq| j	sX|sf| || q| || n\t|}t|tkr| 
|dddd tD t|j|r| |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   s   | ]}|j V  qd S rA   r   r   r+   r+   r,   	<genexpr>  s     z)CodeGenerator.visit_If.<locals>.<genexpr>)r}   testrD   tor
   int1rm   ry   rz   r   r   rN  rD  rL   rF   _condition_typesr!  r$   rQ   r   r~   r   )rf   rS   rA  Zcontains_returnr+   r+   r,   r     s0    
 
 zCodeGenerator.visit_Ifc              
   C   s   |  |j}t|r|jtj| jd}t|  |  \}}| j	 }| j
| tj|  |j| j}| j }| j	 }| j
| tj|  |j| j}| j }| || |j|jkstd|j d|j |j}	|	tjkr|	| jgng }
| j|
|jd}||  |
rH| j|  | j|jg | j|  ||  |
r| j|  | j|jg |
rtj|d|	nd W  5 Q R  S Q R X n^t|}t|tkr|  |d!d"dd	 tD t|j#|r|  |jS |  |jS d S )
Nr  zAternary expression with dynamic condition has inconsistent types r   Tr   rO  rP  c                 s   s   | ]}|j V  qd S rA   r   r   r+   r+   r,   rQ    s     z,CodeGenerator.visit_IfExp.<locals>.<genexpr>)$r}   rR  rD   rS  r
   rT  rm   rd   r   r>  r   r   r   r~   rn   r   r   rF   r'   r   r   rF  r   rG  rH  r  rI  rJ  r   rK  rL   rU  r   r!  r$   rQ   )rf   rS   rA  r   rL  r4  Zthen_valr5  Zelse_valr   Zret_type_irrM  r+   r+   r,   r     sT    




2 zCodeGenerator.visit_IfExpc                 C   s   d S rA   r+   r   r+   r+   r,   
visit_Pass  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	krtt
||kS t|jd tjkrt
||k	S | jt|jd }|d kr| |d|jd j| |||S )Nr   z1simultaneous multiple comparison is not supportedr   z<AST comparison operator '{}' is not (currently) implemented.)r  ZcomparatorsZopsr   r}   r  rL   rF   r   Isr   IsNot_method_name_for_comp_opr   r!  rQ   r  )rf   rS   r  r  	lhs_value	rhs_valuer  r+   r+   r,   visit_Compare  s"     zCodeGenerator.visit_Compare__eq____ne____lt____le____gt____ge__rY  c                 C   s   |  |j}| jt|j}|d kr>| |d|jj dt|rXt	||| j
dS zt	|| W S  tk
r   | |d| dt|j Y nX d S )NzAST unary operator 'z!' is not (currently) implemented.r  z)' is not (currently) implemented on type )r}   operand_method_name_for_unary_opr   rF   r  r   rQ   rD   r   rm   AttributeError)rf   rS   rc  rT   r+   r+   r,   visit_UnaryOp  s     zCodeGenerator.visit_UnaryOp__neg____pos____not__
__invert__rd  c              
      sJ  t }|\}} \}}j }j| j| |j j	  j
}|  g }	g }
g }|D ]}||krvt|| std| dt|| std| d|| j|| jkstd| d|| j d|| j d|	| |
|| j |||  qv|| jfdd	|
D d
d	 |D }j| fdd	|
D  j  t|	D ]8\}}tj ||
| j|< j| j
|< q|j}j  j|j fdd	tt|D  j|  fdd	|
D }j| t|	D ]8\}}tj|||
| j|< j| j
|< q2j| |j j	  j
}g }|D ]}||kr|||  qj!dd	 |D  W 5 Q R X t|	D ]6\}}tj|"||
| }|j|< |j
|< q|j#D ]"}ds4tdt$j%&| q"d S )Nzcannot reassign constxpr z in the loopzcannot reasign constexpr Loop-carried variable  has initial type  but is re-assigned to : in loop! Please make sure that the type stays consistent.c                    s   g | ]}|  jqS r+   rE  r-   rr   r+   r,   r/   J  s     z-CodeGenerator.visit_While.<locals>.<listcomp>c                 S   s   g | ]
}|j qS r+   r   r.   rW   r+   r+   r,   r/   K  s     c                    s   g | ]}|  jqS r+   rE  r-   rr   r+   r,   r/   N  s     c                    s   g | ]}  |qS r+   )rW   r2   )before_blockr+   r,   r/   V  s     c                    s   g | ]}|  jqS r+   rE  r-   rr   r+   r,   r/   Y  s     c                 S   s   g | ]
}|j qS r+   r   r.   yr+   r+   r,   r/   h  s     FzNot implemented)'rd   r   rm   r>  r   r   r   r   r~   poprk   r?  rD   r'   rF   r   Zcreate_while_opZcreate_block_with_parentZ
get_beforerO   r
   r   r   rW   rh   r}   rR  r  Zcreate_condition_opr   ranger  Z	get_afterrI  rK  r   r   r   r   )rf   rS   rB  rj   ro   r   rL  dummyZ	loop_defsr  r   	init_argsr<   Zwhile_opr   rA  Zafter_blockyieldsZnew_defr   r+   )rp  rf   r,   visit_While'  sx    

"
&

 

zCodeGenerator.visit_Whilec                 C   sJ   |j jjdkst| |j}| |j}t|rB|j|| j	dS || S )Nr  r  )
r   	__class__rQ   r'   r}   rK   slicerD   __getitem__rm   )rf   rS   r  Zslicesr+   r+   r,   visit_Subscriptt  s    zCodeGenerator.visit_Subscriptc                    s    fdd|j D S )Nc                    s   g | ]}  |qS r+   r|   )r.   dimrr   r+   r,   r/   }  s     z0CodeGenerator.visit_ExtSlice.<locals>.<listcomp>)Zdimsr   r+   rr   r,   visit_ExtSlice|  s    zCodeGenerator.visit_ExtSlicec                    s    |jj} fdd|jjD }t fdd|jjD }|tjkr|||}t|j	j
|jj
|jj
}|D ]<}t| j|jj<  |j |jD ]}tj | qqnd S d }	|tjkr|||}|j	}
|j}|j}|j}	n|tkrbt|dkr|d n  td}
t|dkr(|d n  |jjd }t|dkrP|d n  td}ntdd	}t|r|j
dk rt|j
 }d
}||
 }
}tj|
 j}
tj| j}tj| j}|
j  r|j  r|j  st!d|
j d|j d|j dtj"#|
j|j}tj"#||j}|$ j}|j%tjjj&j'k}|
j(}
|j(}|j(} j)|
||}
 j)|||} j)|||} j*|} +|jjtj,|| t- }|\}} . \}} j/ } j0|  j12|  |j  j13  |4  g }g }g } j5D ]}||kr.t6 j5| sZt7| dt6|| slt7 j5| j8|| j8kst7d| d|| j8 d j5| j8 d|2| |2tj||  j |2tj j5|  j q. 9||  j:|
||dd |D }|	d k	r8|;d j<|	  j12|  j0|=d |>  _i  _5t?|D ]6\}} +|tj,|=d@|d || j8 qn |j  j13  g } j5D ],}||kr|2tj j5|  j qt|dkr jAdd |D  |=dB }|C dks<t7d j0|=d |D }|rx jE||} jF||
} j|jj j(G|  +|jjtj,|| W 5 Q R X t?|D ],\}} +|tj,|H||| j8 q|jD ]"}d	s t7dtj | qd S )Nc                    s   g | ]}  |qS r+   r|   ro  rr   r+   r,   r/     s     z+CodeGenerator.visit_For.<locals>.<listcomp>c                 3   s   | ]}  |V  qd S rA   r|   r.   keywordrr   r+   r,   rQ    s     z*CodeGenerator.visit_For.<locals>.<genexpr>r   r   r	   zAOnly `range` and `static_range` iterators are currently supportedFTz0For loop bounds and step must all be ints, are (rP  r1  z is not tensorrk  rl  rm  rn  c                 S   s   g | ]
}|j qS r+   r   ro  r+   r+   r,   r/     s     ztt.num_stagesc                 S   s   g | ]
}|j qS r+   r   rq  r+   r+   r,   r/     s     z7We use SCF, so the loop body should only have one blockz)Don't know what to do with else after for)Ir}   iterr   rU   r   keywordsr
   static_rangert  startrK   endstepr   rh   r   r   r   r~   r   r   r   r   
num_stagesr  NumRuntimeErrorrE   r   r   rm   r   r   r   semanticZinteger_promote_implr   r!   r   r    r   Zcreate_int_castZcreate_undefr   r   rd   r   r>  r   r   r   rs  r?  rk   rD   r'   rF   r   Zcreate_for_opZset_attrZget_int32_attrget_bodyri   rO   rW   rI  Z
get_parentsizeZget_induction_varZ
create_subZ
create_addZreplace_all_uses_withrK  )rf   rS   ZIteratorClassZ	iter_argsZiter_kwargsiteratorr  r   r   r  ZlbZubr  Znegative_stepZiv_typeZ
iv_ir_typeZiv_is_signedZivrB  rj   ro   r   rL  blockrv  rw  r  r<   Zfor_opZfor_op_regionr+   rr   r,   	visit_For  s    





&((
$"



$
 

0


 $&
zCodeGenerator.visit_Forc                 C   s0   |  |j}|  |j}|  |j}t|||S rA   )r}   lowerupperr  rz  )rf   rS   r  r  r  r+   r+   r,   visit_Slice  s    zCodeGenerator.visit_Slicec                 C   s   |  |jS rA   )r}   rK   r   r+   r+   r,   visit_Index  s    zCodeGenerator.visit_Indexr{   c                 C   s   |j | |jfS rA   )rW   r}   rK   r   r+   r+   r,   visit_keyword  s    zCodeGenerator.visit_keywordc                 C   sD   | j s
d S | |j}|jd k	r,| |jnd}tjj||| jdS )Nr   r  )r   r}   rR  msgr
   r   device_assertrm   )rf   rS   rR  r  r+   r+   r,   visit_Assert  s
    zCodeGenerator.visit_AssertrT   c                    s  t j|jf |  fdd|jD  dd  D  t }dd t D  fddD }fddt D  dd  D }d	d  D }t|j||}| j	|svt
g |}	|j}
t|\}}|jd kr| jn|j}t| j|	|
||| j||| j|j||| jj| jj|d
}z||  W n8 tk
rb } zt| jj| jd |W 5 d }~X Y nX |j}|| j|< n
| j| }| j|}| j||}|  dks|d krd S |  dkrt!|"d|S g }t#|  D ] }|$t!|"|||  qt%|S d S )Nc                    s   g | ]} | qS r+   r+   )r.   r<   rU   r+   r,   r/     s     z2CodeGenerator.call_JitFunction.<locals>.<listcomp>c                 S   s    g | ]}t |r|nt|qS r+   )rD   r   ro  r+   r+   r,   r/     s     c                 S   s   g | ]\}}t |r|qS r+   )rE   r.   r   rW   r+   r+   r,   r/     s      c                    s   i | ]}| | qS r+   r+   r2   r  r+   r,   r      s      z2CodeGenerator.call_JitFunction.<locals>.<dictcomp>c                    s    g | ]\}}| krd n|qS rA   r+   r  )
constexprsr+   r,   r/   "  s     c                 S   s   g | ]}|d k	r|j qS rA   r   ro  r+   r+   r,   r/   #  s      c                 S   s   g | ]}|d k	r|j qS rA   r   ro  r+   r+   r,   r/   $  s      )
r   r   r   r   rN   r_   ra   r   r   r   r   r   )&r[   getcallargsrT   rR   r   rO   r>   rQ   r   Zhas_functionr
   function_type__globals__rc   r   r   r   r   rN   rm   r   r   r}   r   	Exceptionr   r   rP   r   r   Zget_functioncallZget_num_resultsr   rK  rt  r   rI   )rf   rT   rU   ru   r   r4   arg_vals	arg_typesfn_namer   rz   r_   ra   r   re   eZcallee_ret_typesymbolZcall_opresultsr   r+   )rU   r  r,   call_JitFunction  sV          &
zCodeGenerator.call_JitFunctionc           	   
      sJ  t  |j} j|}|d k	r.| |S t fdd|jD } fdd|jD }|tj	j
krp jspd S t|trt|||  |||S t|drt|jstj	|r&t jd}t|}d|jkrވ |d< z||||W S  tk
r$ } zt jj|d |W 5 d }~X Y nX | j kr@tt |}|||S )Nc                 3   s   | ]}  |V  qd S rA   r|   r  rr   r+   r,   rQ  N  s     z+CodeGenerator.visit_Call.<locals>.<genexpr>c                    s   g | ]}  |qS r+   r|   ro  rr   r+   r,   r/   O  s     z,CodeGenerator.visit_Call.<locals>.<listcomp>__self__r  Z
_generator)rL   r}   r    statically_implemented_functionsr   r   r  rU   r
   r   r  r   rB   r   rX   r  hasattrrD   r  
is_builtinrm   r[   	signature
parametersr  r   r   rP   r   r  r%   )	rf   rS   rT   Zstatic_implementationkwsrU   extra_kwargssigr  r+   rr   r,   r   H  s0    

"

$
zCodeGenerator.visit_Callc                 C   s
   t |jS rA   r   rK   r   r+   r+   r,   visit_Constantj  s    zCodeGenerator.visit_ConstantrS   c                 C   sx   t |jdkr| |d| |jd }| |jd }| jt|j}|d krj| |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   rF   r  r!  rQ   r  r"  r+   r+   r,   visit_BoolOpm  s      zCodeGenerator.visit_BoolOplogical_and
logical_orr  )      c                 C   s
   t |jS rA   r  r   r+   r+   r,   visit_NameConstant}  s    z CodeGenerator.visit_NameConstantc                 C   s
   t |jS rA   )r   r;  r   r+   r+   r,   	visit_Num  s    zCodeGenerator.visit_Numc                 C   s   t t|S rA   )r   r   literal_evalr   r+   r+   r,   	visit_Str  s    zCodeGenerator.visit_Strc                 C   s>   |  |j}t|r2|jdkr2tjj|d| jdS t||jS )NT)r   r   )rm   )	r}   rK   rD   r   r
   r  Zpermuterm   r   )rf   rS   r  r+   r+   r,   r     s
    
zCodeGenerator.visit_Attributec                 C   s   t j| | d S rA   r   r   r+   r+   r,   
visit_Expr  s    zCodeGenerator.visit_Exprc                 C   s   d S rA   r+   r   r+   r+   r,   visit_NoneType  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   )rH   r  rO   rB   r   Constantr"   rK   FormattedValue
conversionr}   rE   r   rF   chrr!  r'   r$   )rf   rS   r  r   rK   Zconversion_codeZ	evaluatedr+   r+   r,   visit_JoinedStr  s"    

*zCodeGenerator.visit_JoinedStrc                    s  |d krd S t   t dt t dt | j}| j }|| _t|dr~t|dr~| j	| j
| j|j |j | j }zt |}W nN tk
r    Y n: tk
r } zt| jj| jt|d W 5 d }~X Y nX |r|| _| j	| |W  5 Q R  S Q R X d S )Nignorelineno
col_offset)warningscatch_warningssimplefilterDeprecationWarningPendingDeprecationWarningr   rm   r   r  r   r_   ra   r  r  superr}   r   r  r   rP   r1   )rf   rS   Z	last_noderL  r=   r  ry  r+   r,   r}     s*    


*zCodeGenerator.visitc                 C   s   |  |dt|jd S )Nzunsupported AST node type: {})r   r!  rF   rQ   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}nHz| |jd }W n2 t	k
r } zdt
| d }W 5 d }~X Y nX 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  rU   r  r   rL   r}   rB   r   NotImplementedErrorr  r1   r   r   rP   )rf   rS   	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                    s@   dd  fdd|j D D } fdd|jD }t||S )Nc                 S   s   i | ]\}}|t |qS r+   )rL   )r.   r<   rK   r+   r+   r,   r     s    z>CodeGenerator.static_executor.<locals>.ret.<locals>.<dictcomp>c                 3   s   | ]}  |V  qd S rA   r|   r  rr   r+   r,   rQ    s     z=CodeGenerator.static_executor.<locals>.ret.<locals>.<genexpr>c                    s   g | ]}t  |qS r+   )rL   r}   ro  rr   r+   r,   r/     s     z>CodeGenerator.static_executor.<locals>.ret.<locals>.<listcomp>)r  rU   r   )rf   rS   r  rU   	python_fnrr   r,   r=     s
    z*CodeGenerator.static_executor.<locals>.ret)r   r   )r  r=   r+   r  r,   static_executor  s    zCodeGenerator.static_executorr  )NNFNFNr   )~rQ   rw   rx   r   r   r   r"   rg   r  rH   rt  floatintrB   r   r   r   r   updater
   r   Zdevice_printminimummaximumr   r   r   r   r   r   r   r   r   r   r   r   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   operatorr9  rD  rN  r   r   rV  r\  EqNotEqLtLtEGtGtErY  cmpoprf  USubUAddNotInvertrd  unaryoprx  r|  r~  r  r  r  r   r  r  r  r   r  BoolOpr  AndOrr  boolopsysversion_infor  r  r  r   r  r  r  r}   r   r   r  r  Zstatic_assertZstatic_printr   r  r   r   __classcell__r+   r+   r  r,   r      s  
            &,
.	C	
            7%5                  M /"&    r   c                 C   sJ   d}t | D ]8\}}|t|7 }||jkr2|d7 }||jkr|d7 }q|S )Nr   r0   d)rO   r"   
equal_to_1divisible_by_16)r  specializationsuffixr   r   r+   r+   r,   kernel_suffix  s    


r  c                    s   j }fdd  fddj D }j }}tj fdd|j	D }	dd |j
D }
| }||	 fddj D }t\}}tg |}t||||||
d	||||d
}|  |j}||_|S )Nc                    s   t | tr j| S | S rA   )rB   r"   rR   index)r   r  r+   r,   <lambda>      zast_to_ttir.<locals>.<lambda>c                    s   i | ]\}} ||qS r+   r+   )r.   keyrK   )cst_keyr+   r,   r      s      zast_to_ttir.<locals>.<dictcomp>c                    s*   i | ]"}|| kr" | d kr"dndqS )i1Tr   r+   r.   k)tysr+   r,   r     s      c                 S   s   i | ]}|d gqS ))ztt.divisibility   r+   r	  r+   r+   r,   r     s      c                    s"   g | ]\}}| j krt|qS r+   )r4   r   )r.   r
  r   )r   r+   r,   r/   
  s     
 zast_to_ttir.<locals>.<listcomp>T)
rz   r4   r   r   r   r   r_   ra   r   r   )attrsr4   itemsr  ri   r1   rH   r  r  r  r  r  rc   r
   r  r   r}   r   r   r   )rT   r   r   r   r   r  r4   rz   r   Znew_constantsZ	new_attrsZall_constantsr  r_   ra   r   re   r=   r+   )r  rT   r   r  r,   ast_to_ttir  s4    


     r  )3r   r[   r  r  r  r   r   typingr   r   r   r   r   r   r   r   r
   Z_C.libtritonr   r   r   r   Zruntime.jitr   runtimer   errorsr   r   r   typesr   r   r>   r   rD   rE   rG   rJ   rL   rX   rc   r  rF   rU  rd   r   ry   r   r  r  r+   r+   r+   r,   <module>   sJ   $
[        #