o
    灛i                     @   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!dede"fddZ#dede"fddZ$dede"fddZ%dede"fddZ&def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 |  r$tjjj}| j|krdnd}|t	| j
 S |  r,t	| S |  rFt| j}dtt	| j}| d| dS |  rLdS J d)	NPiu_SVFzUnsupported type)is_ptr	mangle_ty
element_tyis_intr
   dtype
SIGNEDNESSSIGNEDint_signednessstrint_bitwidthis_floatingis_blockscalarjoinmapshapeis_void)tyr"   prefixeltr+    r0   `/sda-disk/www/egybert/egybert_env/lib/python3.10/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 r0   )r   .0r-   r0   r0   r1   
<listcomp>'       zmangle_fn.<locals>.<listcomp>c                    s"   g | ]}| d t  |  qS )c)reprr3   r   	constantsr0   r1   r4   (      " ._d_'_sq_[]__)r)   sortedreplace)namearg_tysr:   mangled_arg_namesmangled_constantsretr0   r9   r1   	mangle_fn%   s   rJ   oreturnc                 C   
   t | tS N)
isinstancer   rK   r0   r0   r1   _is_triton_tensor1      
rQ   c                 C   rM   rN   )rO   r   rP   r0   r0   r1   _is_constexpr5   rR   rS   c                 C   s    t | o| j  p| jjdkS )Nr   )rQ   typer'   numelrP   r0   r0   r1   _is_triton_scalar9   s    rV   c                 C   s   t | ttfS rN   )rO   listtuplerP   r0   r0   r1   _is_list_like=      rY   c                 C   s   t | tr| jS | S rN   )rO   r   valuerP   r0   r0   r1   _unwrap_if_constexprA      r\   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	enumeraterS   rV   r   src__name__	arg_names)nodefnargsidxargr0   r0   r1   _check_fn_argsE   s   ri   c                 C   sr   | }t |ts|j}t |tr|jjj}t|j\}}t|D ]\}}| 	dr4||7 } ||fS q ||fS )Nzdef )
rO   r   re   __code__co_filenameinspectgetsourcelinesr`   strip
startswith)re   base_fn	file_namelines
begin_linerg   liner0   r0   r1   _get_fn_file_lineO   s   


ru   c                   @   s$   e Zd Zdd Zdd Zdd ZdS )enter_sub_regionc                 C   
   || _ d S rN   )	generator)selfrx   r0   r0   r1   __init__f   rR   zenter_sub_region.__init__c                 C   sL   | j j | _| j j | _i | j _| j j | _| j j	 | _
| j| jfS rN   )rx   lscopecopyliveins
local_defs	prev_defsbuilderget_insertion_blockinsert_blockget_insertion_pointinsert_pointry   r0   r0   r1   	__enter__i   s   zenter_sub_region.__enter__c                 O   s(   | j j| j | j| j _| j| j _d S rN   )rx   r   restore_insertion_pointr   r}   r{   r   r~   )ry   rf   kwargsr0   r0   r1   __exit__r   s   
zenter_sub_region.__exit__N)rb   
__module____qualname__rz   r   r   r0   r0   r0   r1   rv   d   s    	rv   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   rw   rN   )gscope)ry   r   r0   r0   r1   rz   {   rR   zContainsReturnChecker.__init__rL   c                 C   s   |D ]
}|  |r dS qdS )NTFvisit)ry   bodysr0   r0   r1   _visit_stmts~   s
   
z"ContainsReturnChecker._visit_stmtsc                 C   s,   t |tr|js| }t| j|S dS NF)rO   r   r_   parser   r   r   )ry   re   fn_noder0   r0   r1   _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_fieldsrO   rW   ASTr   )ry   rd   rI   r   r[   itemr0   r0   r1   generic_visit   s   
z#ContainsReturnChecker.generic_visitrd   c                 C   sP   t |jtjr"|jj| jv r | j|jj }t||j}| |S dS | 	|jS r   )
rO   r[   r   Nameidr   getattrattrr   r   )ry   rd   r[   re   r0   r0   r1   visit_Attribute   s   
z%ContainsReturnChecker.visit_Attributec                 C   s:   t |jtjkr
dS |j| jv r| j|j }| |S dS r   )rT   ctxr   Storer   r   r   )ry   rd   re   r0   r0   r1   
visit_Name   s   
z ContainsReturnChecker.visit_Namec                 C      dS )NTr0   ry   rd   r0   r0   r1   visit_Return      z"ContainsReturnChecker.visit_Returnc                 C   r   r   r0   r   r0   r0   r1   visit_Assign      z"ContainsReturnChecker.visit_Assignc                 C   r   r   r0   r   r0   r0   r1   visit_AugAssign   r   z%ContainsReturnChecker.visit_AugAssignc                 C      |  |jS rN   r   r   r   r0   r0   r1   visit_Module      z"ContainsReturnChecker.visit_Modulec                 C   r   rN   r   r   r0   r0   r1   visit_FunctionDef   r   z'ContainsReturnChecker.visit_FunctionDefc                 C   s&   |  |j}|jr|p|  |j}|S rN   )r   r   orelse)ry   rd   rI   r0   r0   r1   visit_If   s   zContainsReturnChecker.visit_Ifc                 C   s   |  |jp|  |jS rN   )r   r   r   r   r0   r0   r1   visit_IfExp      z!ContainsReturnChecker.visit_IfExpc                 C   r   rN   )r   funcr   r0   r0   r1   
visit_Call   r   z ContainsReturnChecker.visit_CallN)rb   r   r   rz   boolr   r   r   r   	Attributer   r   r   Returnr   Assignr   	AugAssignr   Moduler   FunctionDefr   Ifr   IfExpr   Callr   r0   r0   r0   r1   r   y   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&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= Z0e1j2d>e1j3d?e1j4d@e1j5dAe1j6dBe1j7dCe1j8dDe1j9dEe1j:dFe1j;dGe1j<dHe1j=dIiZ>ee?e1j@ ef edJ< dKdL ZAdMdN ZBdOdP ZCdQdR ZDdSdT ZEdUdV ZFdWdX ZGe1jHdYe1jIdZe1jJd[e1jKd\e1jLd]e1jMd^iZNee?e1jO ef ed_< d`da ZPe1jQdbe1jRdce1jSdde1jTdeiZUee?e1jV ef edf< dgdh ZWdidj ZXdkdl ZYdmdn ZZdodp Z[dqdr Z\de]eef fdsdtZ^defdudvZ_dwefdxdyZ`dzd{ Zad|d} Zbd~e1jcfddZde1jede1jfdiZgee?e1jh ef ed< eijjdk rdd Zkdd Zldd Zmdd Zndd Zodd Zpdd Zq fddZrdd Zsd~e1jtddf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_typesrq   c                 C   s   || _ t|| _|| _|d | _| j||d || j_|	| j_|d u r*| j n|| _	|d u r3i n|| _
|| _|| _t | _|| _|| _|| _|| _|| _d | _|
d u rY|jn|
| _|| _g | _d | _i | _|  | _d | _d| _d S )Nr   r   F)contextr   r   rq   rs   set_locoptionscodegen_fnscreate_modulemodulefunction_ret_types	prototyper   dictr{   
attributesr:   r   function_name	is_kernelcur_nodedebugr_   	scf_stackret_typer~   _define_name_lookupdereference_namere   visiting_arg_default_value)ry   r   r   r   r   r:   r   r   r   r   r   r   r   r   r_   rq   rs   r0   r0   r1   rz      s4   


zCodeGenerator.__init__c                 C   s   i | ]}|j |qS r0   rb   r3   r   r0   r0   r1   
<dictcomp>   r5   zCodeGenerator.<dictcomp>builtin_namespaceprintminmaxc                 C   s   t | jj||S rN   )r   r   ra   )ry   rd   messager0   r0   r1   _unsupported  s   zCodeGenerator._unsupportedc                 C   sT   t  }| j||}||u rdS t|rdS | jdi | }r(t|dkS dS )NFT__annotations__r   )objectr   getrS   r   )ry   rE   absent_markervalar0   r0   r1   _is_constexpr_global  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 )	NrE   c                    s    j | |S rN   )r{   r   )rE   absentr   r0   r1   local_lookup  s   z7CodeGenerator._define_name_lookup.<locals>.local_lookupc                    s    j | |}||u sA|  jv sAt|tksAt|tsAt|ddsAt|dddsAt|t	j
sA | sA jsAtjdddkrC|S ttd	|  d
dd)N__triton_builtin__Fr    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 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.
 )r   r   r   rT   r   rO   r   r   ro   r
   r    r   r   osenviron	NameErrortextwrapdedentrD   )rE   r   r   r   r0   r1   global_lookup  s0   


	
z8CodeGenerator._define_name_lookup.<locals>.global_lookuprL   c                    s@    }j jfD ]}|| |}||ur|  S q	t|  d)Nz is not defined)r   r   r   )rE   r   lookup_functionr[   r   r   r   ry   r0   r1   name_lookup8  s   
z6CodeGenerator._define_name_lookup.<locals>.name_lookup)r$   r   r   )ry   r   r0   r   r1   r     s
   z!CodeGenerator._define_name_lookuprE   r[   rL   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~   )ry   rE   r[   r0   r0   r1   	set_valueB  s   
zCodeGenerator.set_valuec                 C   s   | j  }| j  }||fS rN   )r   get_locr   )ry   locipr0   r0   r1   _get_insertion_point_and_locK  s   

z*CodeGenerator._get_insertion_point_and_locc                 C   s   | j | | j | d S rN   )r   r   r   )ry   r   r   r0   r0   r1   _set_insertion_point_and_locS  s   z*CodeGenerator._set_insertion_point_and_locc                 C   s8   t |s|g}|D ]}| | t|tjr d S q	d S rN   )rY   r   rO   r   r   )ry   stmtsstmtr0   r0   r1   visit_compound_statementZ  s   
z&CodeGenerator.visit_compound_statementc                 C      t j| | d S rN   r   NodeVisitorr   r   r0   r0   r1   r   f     zCodeGenerator.visit_Modulec                    s0     |j}|d u sJ  fdd|jD }|S )Nc                       g | ]}  |qS r0   r   )r3   r/   r   r0   r1   r4   l      z,CodeGenerator.visit_List.<locals>.<listcomp>)r   r   elts)ry   rd   r   r  r0   r   r1   
visit_Listi  s   zCodeGenerator.visit_Listc                    s     |j}|d u r jg  tj}n8t|tr9 fdd|D }dd |D } jdd |D  t|}ntj	| j} j|j
g |j} jd u rV| _d S  j|krftd j d| d S )Nc                    s   g | ]
}t j| jqS r0   )r
   core
_to_tensorr   r3   vr   r0   r1   r4   z      z.CodeGenerator.visit_Return.<locals>.<listcomp>c                 S      g | ]}|j qS r0   rT   r  r0   r0   r1   r4   {      c                 S   r  r0   handler  r0   r0   r1   r4   |  r  zInconsistent return types:  and )r   r[   r   rI   r
   voidrO   rX   r	  r
  r  rT   r   	TypeError)ry   rd   	ret_valueret_ty
ret_values	ret_typesrI   r0   r   r1   r   p  s"   




zCodeGenerator.visit_Returnc              	   C   s  |  |j\}}| jr| |dt|jjD ]G\}}|jj| d  }|j}|j}tj	|t
 d}	|d u r@tj|	g|d}
ntj|	||d}
z| jrNJ d| _|  |
 W d| _qd| _w | jrddnd	}| j| j| j| j| j|| j| _| j| j | j }g }d
}t|D ]J\}}|| jv r| j| }t|st| j| }|| q|| jv r| j| D ]\}}| j||| q|t| j|| jj|  |d7 }q| j  }t!||D ]
\}}| "|| q| j#| | $|j% | j&d u s| j&t'j(krt'j(| _&| j)g  n+t*| j&t+r+t,| j&| j_-| j.| j| j n| j&g| j_-| j.| j| j |rE| j/| | j0  d S )Nz,nested function definition is not supported.r   r   r   targetsr[   )targetr[   
annotationTFpublicprivater   )1r   rf   re   r   r`   defaultsr  rh   r   r   r   r   	AnnAssignr   r   r   get_or_insert_functionr   r   r   to_irr_   	push_backadd_entry_blockr:   rS   r   appendr   set_arg_attrr   param_typesr   zipr   set_insertion_point_to_startr   r   r   r
   r  rI   rO   rX   rW   r  
reset_typeset_insertion_point_to_endfinalize)ry   rd   rc   kwarg_namesr   default_valuearg_noder  rE   	st_target	init_node
visibilityentry
arg_valuesrg   arg_namecstr[   	insert_pt	arg_valuer0   r0   r1   r     sh   





 

zCodeGenerator.visit_FunctionDefc                 C   s4   g }|j D ]
}|| |g7 }q| |j}||fS rN   )rf   r   kwarg)ry   rd   rc   rh   r/  r0   r0   r1   visit_arguments  s
   
zCodeGenerator.visit_argumentsc                 C   s   t j| | |jS rN   )r   r  r   rh   r   r0   r0   r1   	visit_arg  s   zCodeGenerator.visit_argc                 C   sr   |  |j}|  |j}|  |j}|tkr4|| jv r"t| dt|s*t|}|| j|< | j| S | |S )Nz4 is already defined. constexpr cannot be reassigned.)	r   r  r  r[   r   r{   
ValueErrorrS   r   )ry   rd   r  r  r[   r0   r0   r1   visit_AnnAssign  s   



zCodeGenerator.visit_AnnAssignc           	      C   s   g }|j D ]
}|| |g7 }qt|dkr| |d|d }| |j}t|s-|g}t|s4|g}tjf}t||D ]#\}}t	|}|d urZt
|sZt||sZtj|| j}| || q=d S )Nr   z2simultaneous multiple assignment is not supported.r   )r  r   lenr   r[   rY   r
   r    r*  r\   rQ   rO   r	  r
  r   r   )	ry   rd   _namesr  namesvaluesnative_nontensor_typesrE   r[   r0   r0   r1   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opr[   r   r   r   )ry   rd   rE   lhsrhsassignr0   r0   r1   r     s   

zCodeGenerator.visit_AugAssignc                 C   s"   t |jtjkr|jS | |jS rN   )rT   r   r   r   r   r   r   r0   r0   r1   r     s   zCodeGenerator.visit_Namec                 C   r  rN   r  r   r0   r0   r1   visit_Store  r  zCodeGenerator.visit_Storec                 C   r  rN   r  r   r0   r0   r1   
visit_Load  r  zCodeGenerator.visit_Loadc                    s    fdd|j D }t|S )Nc                    r  r0   r   )r3   xr   r0   r1   r4     r  z-CodeGenerator.visit_Tuple.<locals>.<listcomp>)r  rX   )ry   rd   rf   r0   r   r1   visit_Tuple  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__)rQ   r   r   resub)ry   method_namerH  rI  reverse_method_namer0   r0   r1   _apply_binary_method  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   rT   rG  r   formatrb   rU  ry   rd   rH  rI  rS  r0   r0   r1   visit_BinOp  s   zCodeGenerator.visit_BinOp__add____sub____mul____truediv____floordiv____mod____pow__
__lshift__
__rshift____and____or____xor__rX  c                 C   s  | j | | |j | j  }| j }i }|jr9| j | | | _i | _| |j | j }| j  }g }g }g }	|D ]}
|df|dffD ](\}}|
|v rs||
 j	||
 j	kssJ d|
 d||
 j	 d| d||
 j	 qK|
|v s||
|v r|
|
 |
|
|v r||
 j	n||
 j	 |	
|
|v r||
 j n||
 j  |
|v r|
|vr||
 ||
< |
|v r|
|vr||
 ||
< qA| | @ D ]7}
|
|v rq||
 j	}||
 j	}||ksJ d|
 d| d	| d
|
|
 |
| |	
||
 j  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   r+  r   r   r   r~   r|   r   r{   rT   r'  r  get_typekeys)ry   rd   r}   
then_block
else_block	then_defs	else_defsrB  r  ir_ret_typesrE   defs
block_namethen_tyelse_tyr0   r0   r1   visit_then_else_blocks7  sj   





"




z$CodeGenerator.visit_then_else_blocksc                    sx  d}t | }|\}}| j }| j }| j }	| j| | j|j|| | ||||\ }}}
}}| j| | rL| rLd}|	  |	 s`|r`| j
|	fdd|
D  | j| |	 sz|rz| j
|	 fdd|
D  |r|D ]}|	| q~W d    n1 sw   Y  |r| j|	 t|
D ]\}}tj|	||| }| || qd S d S )NTFc                       g | ]} | j qS r0   r  r3   nro  r0   r1   r4     r  z4CodeGenerator.visit_if_top_level.<locals>.<listcomp>c                    rw  r0   r  rx  rp  r0   r1   r4     r  )rv   r   create_blockr-  create_cond_branchr  rv  
has_returnerasehas_terminatorcreate_branchadd_argumentr+  r`   r
   r	  r   rh   r   )ry   condrd   has_endif_blocksrr}   ip_blockrm  rn  endif_blockrB  r  rq  r-   r   rE   
new_tensorr0   )rp  ro  r1   visit_if_top_leveln  s>   



z CodeGenerator.visit_if_top_levelc                    sv  t }|\}} \}}j }|jrj nd }	||||	\ }}	}
}}|| jfdd|D |jd}|	|
  j|
  t|
dkrejfdd|
D  |jsm| }	n|		|  j|  t|
dkrj fdd|
D  W d    n1 sw   Y  t|
D ]\}}tj|||| }|| qd S )Nc                       g | ]}|  jqS r0   r$  r   r2   r   r0   r1   r4         z.CodeGenerator.visit_if_scf.<locals>.<listcomp>Tr   c                    rw  r0   r  rx  rz  r0   r1   r4     r  c                    rw  r0   r  rx  r{  r0   r1   r4     r  )rv   r   r   r|  r   rv  r   create_if_opr  merge_block_beforeget_then_blockr-  r@  create_yield_opget_else_blockr`   r
   r	  r   
get_resultr   )ry   r  rd   r  r}   r   r   last_locrm  rn  rB  r  if_opr   rE   r  r0   )rp  ry   ro  r1   visit_if_scf  s2   

 
zCodeGenerator.visit_if_scfc              	   C   s   |  |j}t|r;|jtj| jd}t| j |}| j	r&|r&| 
|d| j	s+|s3| || d S | || d S t|}t|tvrZ| 
|dddd tD t|j|rd| |j d S | |j d S )NrO  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 rN   r   r   r0   r0   r1   	<genexpr>      z)CodeGenerator.visit_If.<locals>.<genexpr>)r   testrQ   tor
   int1r   r   r   r   r   r  r  r\   rT   _condition_typesrY  r)   rb   r   r   r   )ry   rd   r  contains_returnr0   r0   r1   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	 }| j
| tj|  |j| j}| j }| || |j|jksnJ 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 )
NrO  zAternary expression with dynamic condition has inconsistent types r  Tr   r  r  c                 s   r  rN   r   r   r0   r0   r1   r    r  z,CodeGenerator.visit_IfExp.<locals>.<genexpr>)#r   r  rQ   r  r
   r  r   rv   r   r|  r+  r	  r
  r   r   r   r   rT   r  r$  r  r  r  r  r-  r  r  r   r  r\   r  r   rY  r)   rb   )ry   rd   r  r   r  rm  then_valrn  else_valr   ret_type_irr  r0   r0   r1   r     sT   




$#zCodeGenerator.visit_IfExpc                 C      d S rN   r0   r   r0   r0   r1   
visit_Pass  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	kr:t
||u S t|jd tjk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   rV  r\   rT   r   Isr   IsNot_method_name_for_comp_opr   rY  rb   rU  )ry   rd   rH  rI  	lhs_value	rhs_valuerS  r0   r0   r1   visit_Compare  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.rO  z)' is not (currently) implemented on type )r   operand_method_name_for_unary_opr   rT   rG  r   rb   rQ   r   r   AttributeError)ry   rd   r  re   r0   r0   r1   visit_UnaryOp  s   zCodeGenerator.visit_UnaryOp__neg____pos____not__
__invert__r  c              
      sF  t j}|\}} \}}j }j| j| |j j	  j
}|  g }	g }
g }|D ]T}||v rt|| sOJ d| dt|| s]J d| d|| j|| jks{J d| d|| j d|| j d|	| |
|| j |||  q;|| jfdd	|
D d
d	 |D }j| fdd	|
D  j  t|	D ]\}}tj ||
| j|< j| j
|< qÈ|j}j  j|j fdd	tt|D  j| fdd	|
D }j| t|	D ]\}}tj|||
| j|< j| j
|< qj| |j j	  j
}g }|D ]}||v r[|||  qMj dd	 |D  W d    n	1 ssw   Y  t|	D ]\}}tj|!||
| }|j|< |j
|< q||j"D ]}J d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                    r  r0   r  r2   r   r0   r1   r4   J  r  z-CodeGenerator.visit_While.<locals>.<listcomp>c                 S   r  r0   r  r3   rh   r0   r0   r1   r4   K  r  c                    r  r0   r  r2   r   r0   r1   r4   N  r  c                    r  r0   )rh   r8   )before_blockr0   r1   r4   V  r  c                    r  r0   r  r2   r   r0   r1   r4   Y  r  c                 S   r  r0   r  r3   yr0   r0   r1   r4   h  r  FzNot implemented)&rv   r   r   r|  r+  r   r'  r   r   popr~   r  rQ   rT   r   create_while_opcreate_block_with_parent
get_beforer`   r
   r	  r   rh   r{   r   r  r-  create_condition_opr  ranger@  	get_afterr  r  r   r   r  r   )ry   rd   r  r}   r   r   r  dummy	loop_defsrB  r  	init_argsrE   while_opr   r  after_blockyieldsnew_defr   r0   )r  ry   r1   visit_While'  s   


&

C

zCodeGenerator.visit_Whilec                 C   sJ   |j jjdks	J | |j}| |j}t|r!|j|| jdS || S )NrE  rO  )	r   	__class__rb   r   r[   slicerQ   __getitem__r   )ry   rd   rH  slicesr0   r0   r1   visit_Subscriptt  s   zCodeGenerator.visit_Subscriptc                    s    fdd|j D S )Nc                    r  r0   r   )r3   dimr   r0   r1   r4   }  r  z0CodeGenerator.visit_ExtSlice.<locals>.<listcomp>)dimsr   r0   r   r1   visit_ExtSlice|  r]   zCodeGenerator.visit_ExtSlicec                    s    |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 }	|tju ru||i |}|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|
 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- x}|\}} . \}} j/ } j0|  j12|  |j  j13  |4  g }g }g } j5D ]a}||v rt6 j5| sJ | dt6|| sJ  j5| j7|| j7ksJ d| d|| j7 d j5| j7 d|2| |2tj||  j |2tj j5|  j q 8||  j9|
||dd |D }|	d ur|:d j;|	  j12|  j0|<d |=  _i  _5t>|D ]\}} +|tj,|<d?|d || j7 q3 |j  j13  g } j5D ]}||v rt|2tj j5|  j q_t|dkr j@dd |D  |<dA }|B dksJ d j0|<d |C }|r jD||} jE||
} j|jj j(F|  +|jjtj,|| W d    n	1 sw   Y  t>|D ]\}} +|tj,|G||| j7 q|jD ]}J dd S )Nc                    r  r0   r   r  r   r0   r1   r4     r  z+CodeGenerator.visit_For.<locals>.<listcomp>c                 3       | ]}  |V  qd S rN   r   r3   keywordr   r0   r1   r        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 (r  rj  z is not tensorr  r  r  r  c                 S   r  r0   r  r  r0   r0   r1   r4     r  ztt.num_stagesc                 S   r  r0   r  r  r0   r0   r1   r4     r  z7We use SCF, so the loop body should only have one blockz)Don't know what to do with else after for)Hr   iterr   rf   r   keywordsr
   static_ranger  startr[   endstepr   r{   r  r   r   r   r   r   r  r   
num_stagesr@  NumRuntimeErrorrS   r	  r
  r   r    r   r  semanticinteger_promote_implr$  r#   r!   r"   r  create_int_castcreate_undefr   r   rv   r   r|  r+  r   r'  r  r  r~   rQ   rT   r   create_for_opset_attrget_int32_attrget_bodyr|   r`   rh   r  
get_parentsizeget_induction_var
create_sub
create_addreplace_all_uses_withr  )ry   rd   IteratorClass	iter_argsiter_kwargsiteratorr  r   r   r  lbubr  negative_stepiv_type
iv_ir_typeiv_is_signedivr  r}   r   r   r  blockr  r  rB  rE   for_opfor_op_regionr0   r   r1   	visit_For  s   


$&&
 "







0


C&
zCodeGenerator.visit_Forc                 C   s0   |  |j}|  |j}|  |j}t|||S rN   )r   lowerupperr  r  )ry   rd   r  r  r  r0   r0   r1   visit_Slice  s   zCodeGenerator.visit_Slicec                 C   r   rN   )r   r[   r   r0   r0   r1   visit_Index  r   zCodeGenerator.visit_Indexc                 C   s   |j | |jfS rN   )rh   r   r[   r   r0   r0   r1   visit_keyword  r  zCodeGenerator.visit_keywordc                 C   sD   | j sd S | |j}|jd ur| |jnd}tjj||| jdS )Nr   rO  )r   r   r  msgr
   r	  device_assertr   )ry   rd   r  r	  r0   r0   r1   visit_Assert  s
   zCodeGenerator.visit_Assertre   c                    s  t j|jg R i |  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	|st
g |}	|j}
t|\}}|jd u rp| jn|j}t| 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|}| j||}|  dks|d u rd S |  dkrt!|"d|S g }t#|  D ]}|$t!|"|||  qt%|S )Nc                    s   g | ]} | qS r0   r0   )r3   rE   rf   r0   r1   r4     r5   z2CodeGenerator.call_JitFunction.<locals>.<listcomp>c                 S   s    g | ]}t |r
|nt|qS r0   )rQ   r   r  r0   r0   r1   r4          c                 S   s   g | ]
\}}t |r|qS r0   )rS   r3   r   rh   r0   r0   r1   r4     r  c                    s   i | ]}| | qS r0   r0   r8   r  r0   r1   r      r  z2CodeGenerator.call_JitFunction.<locals>.<dictcomp>c                    s    g | ]\}}| v rd n|qS rN   r0   r  )
constexprsr0   r1   r4   "  r  c                 S      g | ]	}|d ur|j qS rN   r  r  r0   r0   r1   r4   #      c                 S   r  rN   r  r  r0   r0   r1   r4   $  r  )
r   r   r   r   r_   rq   rs   r   r   r   r   r   )&rl   getcallargsre   rc   r   r`   rJ   rb   r   has_functionr
   function_type__globals__ru   r   r   r   r   r_   r   r   r   r   r   	Exceptionr   r   ra   r   r   get_functioncallget_num_resultsr   r  r  r'  rX   )ry   re   rf   r   r   r:   arg_vals	arg_typesfn_namer   r   rq   rs   r   rx   ecallee_ret_typesymbolcall_opresultsr   r0   )rf   r  r1   call_JitFunction  sN   
zCodeGenerator.call_JitFunctionc           	   
      sB  t  |j} j|}|d ur| |S t fdd|jD } fdd|jD }|tj	j
u r8 js8d S t|trJt|||  |||S t|drTt|jsZtj	|rt jd}t|}d|jv rn |d< z
||i ||W S  ty } z	t jj|d |d }~ww | j v rtt |}||i |S )Nc                 3   r  rN   r   r  r   r0   r1   r  N  r  z+CodeGenerator.visit_Call.<locals>.<genexpr>c                    r  r0   r   r  r   r0   r1   r4   O  r  z,CodeGenerator.visit_Call.<locals>.<listcomp>__self__rO  
_generator)r\   r   r    statically_implemented_functionsr   r   r  rf   r
   r	  r
  r   rO   r   ri   r"  hasattrrQ   r#  
is_builtinr   rl   	signature
parametersr  r   r   ra   r   rC  r*   )	ry   rd   re   static_implementationkwsrf   extra_kwargssigr  r0   r   r1   r   H  s4   

 

	
zCodeGenerator.visit_Callc                 C   
   t |jS rN   r   r[   r   r0   r0   r1   visit_Constantj  rR   zCodeGenerator.visit_Constantrd   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@  rC  r   r   _method_name_for_bool_opr   rT   rG  rY  rb   rU  rZ  r0   r0   r1   visit_BoolOpm  s   zCodeGenerator.visit_BoolOplogical_and
logical_orr1  )      c                 C   r.  rN   r/  r   r0   r0   r1   visit_NameConstant}  rR   z CodeGenerator.visit_NameConstantc                 C   r.  rN   )r   ry  r   r0   r0   r1   	visit_Num  rR   zCodeGenerator.visit_Numc                 C   s   t t|S rN   )r   r   literal_evalr   r0   r0   r1   	visit_Str  rZ   zCodeGenerator.visit_Strc                 C   s>   |  |j}t|r|jdkrtjj|d| jdS t||jS )NT)r   r   )r   )	r   r[   rQ   r   r
   r  permuter   r   )ry   rd   rH  r0   r0   r1   r     s
   
zCodeGenerator.visit_Attributec                 C   r  rN   r  r   r0   r0   r1   
visit_Expr  r  zCodeGenerator.visit_Exprc                 C   r  rN   r0   r   r0   r0   r1   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   )rW   rC  r`   rO   r   Constantr$   r[   FormattedValue
conversionr   rS   r   rT   chrrY  AssertionErrorr)   )ry   rd   rC  r   r[   conversion_code	evaluatedr0   r0   r1   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   rq   rs   rI  rJ  superr   r   r  r   ra   r7   )ry   rd   	last_noder  rI   r  r  r0   r1   r     s0   


$zCodeGenerator.visitc                 C   s   |  |dt|j)Nzunsupported AST node type: {})r   rY  rT   rb   r   r0   r0   r1   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@  rf   r  r  r\   r   rO   r   NotImplementedErrorr  r7   r   r   ra   )ry   rd   	arg_countpassedr   r  r0   r0   r1   execute_static_assert  s*   


z#CodeGenerator.execute_static_assertc                    s   dt jf fdd}|S )Nrd   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 r0   )r\   )r3   rE   r[   r0   r0   r1   r     s    z>CodeGenerator.static_executor.<locals>.ret.<locals>.<dictcomp>c                 3   r  rN   r   r  r   r0   r1   r    r  z=CodeGenerator.static_executor.<locals>.ret.<locals>.<genexpr>c                    s   g | ]	}t  |qS r0   )r\   r   r  r   r0   r1   r4     r  z>CodeGenerator.static_executor.<locals>.ret.<locals>.<listcomp>)r  rf   r   )ry   rd   r+  rf   	python_fnr   r1   rI     s
   z*CodeGenerator.static_executor.<locals>.ret)r   r   )rY  rI   r0   rX  r1   static_executor  s   zCodeGenerator.static_executorr%  )NNFNFNr   )~rb   r   r   r   r   r   r$   rz   r@  rW   r  floatintrO   r   r   r   r   updater
   r	  device_printminimummaximumr   r   r   r   r   r   r   r   r   r   r   r  r   r   r<  r=  r?  r   r   r   rK  rL  rN  rU  r[  r   AddSubMultDivFloorDivModPowLShiftRShiftBitAndBitOrBitXorrX  r   operatorrv  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   r0  BoolOpr2  AndOrr1  boolopsysversion_infor7  r8  r:  r   r=  r>  rG  r   r   r   rW  rZ  static_assertstatic_printr   r%  r   r   __classcell__r0   r0   rR  r1   r      s   
 
,&
.	C	
7%5$M /"&(r   c                 C   sJ   d}t | D ]\}}|t|7 }||jv r|d7 }||jv r"|d7 }q|S )Nr   r6   d)r`   r$   
equal_to_1divisible_by_16)r(  specializationsuffixr   r   r0   r0   r1   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 rN   )rO   r$   rc   index)r   )re   r0   r1   <lambda>  r  zast_to_ttir.<locals>.<lambda>c                    s   i | ]	\}} ||qS r0   r0   )r3   keyr[   )cst_keyr0   r1   r      r  zast_to_ttir.<locals>.<dictcomp>c                    s*   i | ]}|| v r | d krdndqS )i1Tr   r0   r3   k)tysr0   r1   r     s   * c                 S   s   i | ]}|d gqS ))ztt.divisibility   r0   r  r0   r0   r1   r     r5   c                    s"   g | ]\}}| j vrt|qS r0   )r:   r   )r3   r  r  )r  r0   r1   r4   
  r;   zast_to_ttir.<locals>.<listcomp>T)
r   r:   r   r   r   r   rq   rs   r   r   )attrsr:   itemsr  r|   r7   rW   r(  rC  r  r  r]  ru   r
   r  r   r   r   r   r   )re   r  r   r   r   r  r:   r   r   new_constants	new_attrsall_constantsr  rq   rs   r   rx   rI   r0   )r  re   r  r  r1   ast_to_ttir  s*   


r  )3r   rl   rQ  r~  rK  r   r   typingr   r   r   r   r   r   r   r   r
   _C.libtritonr   r   r   r   runtime.jitr   runtimer   errorsr   r   r   typesr   r   rJ   r   rQ   rS   rV   rY   r\   ri   ru   r\  rT   r  rv   r  r   r   r  r  r0   r0   r0   r1   <module>   sL    $
[        #