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mZmZ ddlmZmZ dd	lmZ d
dlm Z m!Z!m"Z" d dl#m$Z$ 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de'fddZ,dd Z-e'e.e/dhZ0G dd  d Z1G d!d" d"e j2Z3G d#d$ d$e j2Z4d%d& Z5d'd( Z6dS ))    N)AnyCallableDictOptionalTupleTypeUnion   )language)ir)	constexprtensor	str_to_ty)_unwrap_if_constexprnv_tma_desc_type_value)_normalize_tyget_jit_fn_file_line)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 td|  )NPiu_SVz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	TypeError)tyr&   prefixeltr/    r5   \/home/ubuntu/transcripts/venv/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 r5   )r!   .0r2   r5   r5   r6   
<listcomp>(       zmangle_fn.<locals>.<listcomp>c                    s"   g | ]}| d t  |  qS )c)reprr8   r   	constantsr5   r6   r9   )      " ._d_'_sq_[]__)r-   sortedreplace)namearg_tysr?   mangled_arg_namesmangled_constantsretr5   r>   r6   	mangle_fn&   s   rO   oreturnc                 C   
   t | tS N)
isinstancer   rP   r5   r5   r6   _is_triton_value2      
rV   c                 C   rR   rS   )rT   r   rU   r5   r5   r6   _is_triton_tensor6   rW   rX   c                 C   rR   rS   )rT   r   rU   r5   r5   r6   _is_constexpr:   rW   rY   c                 C   s    t | o| j  p| jjdkS )Nr   )rX   typer+   numelrU   r5   r5   r6   _is_triton_scalar>   s    r\   c                 C   s   t | ttfS rS   )rT   listtuplerU   r5   r5   r6   _is_list_likeB      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	enumeraterY   r\   r   src__name__	arg_names)nodefnargsidxargr5   r5   r6   _check_fn_argsF   s   rl   c                   @   s$   e Zd Zdd Zdd Zdd ZdS )enter_sub_regionc                 C   
   || _ d S rS   )	generator)selfro   r5   r5   r6   __init__U   rW   zenter_sub_region.__init__c                 C   sL   | j j | _| j j | _i | j _| j j | _| j j	 | _
| j| jfS rS   )ro   lscopecopyliveins
local_defs	prev_defsbuilderget_insertion_blockinsert_blockget_insertion_pointinsert_pointrp   r5   r5   r6   	__enter__X   s   zenter_sub_region.__enter__c                 O   s(   | j j| j | j| j _| j| j _d S rS   )ro   rw   restore_insertion_pointr{   rt   rr   rv   ru   )rp   ri   kwargsr5   r5   r6   __exit__a   s   
zenter_sub_region.__exit__N)re   
__module____qualname__rq   r}   r   r5   r5   r5   r6   rm   S   s    	rm   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   rn   rS   )gscope)rp   r   r5   r5   r6   rq   j   rW   zContainsReturnChecker.__init__rQ   c                    s   t  fdd|D S )Nc                 3       | ]}  |V  qd S rS   visit)r8   sr|   r5   r6   	<genexpr>n       z5ContainsReturnChecker._visit_stmts.<locals>.<genexpr>)any)rp   bodyr5   r|   r6   _visit_stmtsm   s   z"ContainsReturnChecker._visit_stmtsc                 C   s,   t |tr|js| }t| j|S dS NF)rT   r   rb   parser   r   r   )rp   rh   fn_noder5   r5   r6   _visit_functionp   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_fieldsrT   r]   ASTr   )rp   rg   rN   r   valueitemr5   r5   r6   generic_visitw   s   
z#ContainsReturnChecker.generic_visitrg   c                 C   sP   t |jtjr"|jj| jv r | j|jj }t||j}| |S dS | 	|jS r   )
rT   r   r   Nameidr   getattrattrr   r   )rp   rg   r   rh   r5   r5   r6   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   )rZ   ctxr   Storer   r   r   )rp   rg   rh   r5   r5   r6   
visit_Name   s   
z ContainsReturnChecker.visit_Namec                 C      dS )NTr5   rp   rg   r5   r5   r6   visit_Return      z"ContainsReturnChecker.visit_Returnc                 C   r   r   r5   r   r5   r5   r6   visit_Assign      z"ContainsReturnChecker.visit_Assignc                 C   r   r   r5   r   r5   r5   r6   visit_AugAssign   r   z%ContainsReturnChecker.visit_AugAssignc                 C      |  |jS rS   r   r   r   r5   r5   r6   visit_Module      z"ContainsReturnChecker.visit_Modulec                 C   r   rS   r   r   r5   r5   r6   visit_FunctionDef   r   z'ContainsReturnChecker.visit_FunctionDefc                 C   s&   |  |j}|jr|p|  |j}|S rS   )r   r   orelse)rp   rg   rN   r5   r5   r6   visit_If   s   zContainsReturnChecker.visit_Ifc                 C   s   |  |jp|  |jS rS   )r   r   r   r   r5   r5   r6   visit_IfExp      z!ContainsReturnChecker.visit_IfExpc                 C   r   rS   )r   funcr   r5   r5   r6   
visit_Call   r   z ContainsReturnChecker.visit_CallN)re   r   r   rq   boolr   r   r   r   	Attributer   r   r   Returnr   Assignr   	AugAssignr   Moduler   FunctionDefr   Ifr   IfExpr   Callr   r5   r5   r5   r6   r   h   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\dsdt Z]de^eef fdudvZ_defdwdxZ`dyefdzd{Zad|d} Zbd~d Zcde1jdfddZee1jfde1jgdiZhee?e1ji ef ed< ejjkdk rdd Zldd Zmdd Zndd Zodd Zpdd Zqdd Zr fddZsdd Ztde1juddfddZvdd ZwejjxevejjyewezeeweeeweiZ{ee|e}e1jugef f ed<   Z~S )CodeGeneratorNFr   jit_fnfunction_types	file_namec                 C   sP  || _ 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 | _|| _|| _|| _|| _|| _d | _|| _g | _d | _i | _|  | _ d | _!d| _"d S )Nr   r   r    F)#contextr   rw   r   
begin_lineset_locoptionscodegen_fns
module_mapcreate_modulemodulefunction_ret_types	prototyper   itemsrT   r   getre   r   rr   
attributesr?   r   function_name	is_kernelcur_noderb   	scf_stackret_typeru   _define_name_lookupdereference_namerh   visiting_arg_default_value)rp   r   r   r   r   r?   r   r   r   r   r   r   r   r   rb   r   r   kvmodule_namer5   r5   r6   rq      sD   



zCodeGenerator.__init__c                 C   s   i | ]}|j |qS r5   re   r8   r   r5   r5   r6   
<dictcomp>   r:   zCodeGenerator.<dictcomp>builtin_namespaceprintminmaxc                 C   s   t | jj||S rS   )r   r   rd   )rp   rg   messager5   r5   r6   _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   r   rY   r   )rp   rJ   absent_markervalar5   r5   r6   _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 )	NrJ   c                    s    j | |S rS   )rr   r   )rJ   absentr|   r5   r6   local_lookup  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 |  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 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   r   rZ   r   rT   r   r   
startswithr
   r$   r   r   osenviron	NameErrortextwrapdedentrI   )rJ   r   r   r|   r5   r6   global_lookup  s&   


z8CodeGenerator._define_name_lookup.<locals>.global_lookuprQ   c                    s@    }j jfD ]}|| |}||ur|  S q	t|  d)Nz is not defined)r   r   r   )rJ   r   lookup_functionr   r   r   r   rp   r5   r6   name_lookup1  s   
z6CodeGenerator._define_name_lookup.<locals>.name_lookup)r(   r   r   )rp   r   r5   r   r6   r     s
   z!CodeGenerator._define_name_lookuprJ   r   rQ   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)rr   ru   )rp   rJ   r   r5   r5   r6   	set_value;  s   
zCodeGenerator.set_valuec                 C   s   | j  }| j  }||fS rS   )rw   get_locrz   )rp   locipr5   r5   r6   _get_insertion_point_and_locD  s   

z*CodeGenerator._get_insertion_point_and_locc                 C   s   | j | | j | d S rS   )rw   r~   r   )rp   r   r   r5   r5   r6   _set_insertion_point_and_locL  s   z*CodeGenerator._set_insertion_point_and_locc                 C   s8   t |s|g}|D ]}| | t|tjr d S q	d S rS   )r_   r   rT   r   r   )rp   stmtsstmtr5   r5   r6   visit_compound_statementS  s   
z&CodeGenerator.visit_compound_statementc                 C      t j| | d S rS   r   NodeVisitorr   r   r5   r5   r6   r   _     zCodeGenerator.visit_Modulec                    s0     |j}|d u sJ  fdd|jD }|S )Nc                       g | ]}  |qS r5   r   )r8   r4   r|   r5   r6   r9   e      z,CodeGenerator.visit_List.<locals>.<listcomp>)r   r   elts)rp   rg   r   r	  r5   r|   r6   
visit_Listb  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 rU| _n j|kretd j d|  j } j| d S )Nc                    s   g | ]
}t j| jqS r5   )r
   semantic	to_tensorrw   r8   r   r|   r5   r6   r9   o      z.CodeGenerator.visit_Return.<locals>.<listcomp>c                 S      g | ]}|j qS r5   rZ   r  r5   r5   r6   r9   p      c                 S   r  r5   handler  r5   r5   r6   r9   q  r  zInconsistent return types:  and )r   r   rw   rN   r
   voidrT   r^   r  r  r  rZ   r   r1   create_blockset_insertion_point_to_end)rp   rg   	ret_valueret_ty
ret_values	ret_typesrN   post_ret_blockr5   r|   r6   r   i  s$   




zCodeGenerator.visit_Returnc              	      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 }g }d}tt|D ]Y}| jv r j| }t|st j| }|| q| jv rȈ j| D ]\}} j||| qt  jj!| t"rو j|dd |t# j| jj!|  |d7 }q j$ }t%||D ]
\}} &|| q j'|  (|j)  j$ * rJ  j+d u s& j+t,j-kr1t,j- _+ j.g  n,t  j+t/r=t0 j+n j+g j_1 j2 j j  j. fdd jj1D   j3  |rm j4| d S d S )Nz,nested function definition is not supported.r   r   r   targetsr   )targetr   
annotationTFpublicprivater   ztt.nv_tma_descc                    s*   g | ]} j d ur j| jqS rS   )r   rw   create_poisonto_irr7   r|   r5   r6   r9     s
    
z3CodeGenerator.visit_FunctionDef.<locals>.<listcomp>)5r   ri   rh   r   rc   defaultsr"  rk   r   r   r   r   	AnnAssignr   r   rw   get_or_insert_functionr   r   r   r&  rb   	push_backadd_entry_blockrangelenr?   rY   r   appendr   set_arg_attrrT   param_typesr   r   rx   zipr   set_insertion_point_to_startr  r   has_terminatorr   r
   r  rN   r^   r]   r  
reset_typefinalizer  )rp   rg   rf   kwarg_namesr   default_valuearg_noder"  rJ   	st_target	init_node
visibilityentry
arg_valuesrj   cstr   	insert_ptarg_name	arg_valuer5   r|   r6   r     sp   





 

$
zCodeGenerator.visit_FunctionDefc                 C   s4   g }|j D ]
}|| |g7 }q| |j}||fS rS   )ri   r   kwarg)rp   rg   rf   rk   r6  r5   r5   r6   visit_arguments  s
   
zCodeGenerator.visit_argumentsc                 C   s   t j| | |jS rS   )r   r  r   rk   r   r5   r5   r6   	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   rr   
ValueErrorrY   r   )rp   rg   r"  r!  r   r5   r5   r6   visit_AnnAssign  s   



zCodeGenerator.visit_AnnAssignc           	      C   s   g }t |tjr|| |jg7 }n|jD ]
}|| |g7 }qt|dkr,| |d|d }| |j}t	|s=|g}t	|sD|g}t
jf}t||D ]#\}}t|}|d urjt|sjt ||sjt
j|| j}| || qMd S )Nr   z2simultaneous multiple assignment is not supported.r   )rT   r   r(  r   r!  r   r-  r   r   r_   r
   r$   r1  r   rV   r  r  rw   r   )	rp   rg   _namesr!  namesvaluesnative_nontensor_typesrJ   r   r5   r5   r6   r     s0   
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   )rp   rg   rJ   lhsrhsassignr5   r5   r6   r     s   

zCodeGenerator.visit_AugAssignc                 C   s"   t |jtju r|jS | |jS rS   )rZ   r   r   r   r   r   r   r5   r5   r6   r   	  s   zCodeGenerator.visit_Namec                 C   r  rS   r  r   r5   r5   r6   visit_Store  r  zCodeGenerator.visit_Storec                 C   r  rS   r  r   r5   r5   r6   
visit_Load  r  zCodeGenerator.visit_Loadc                    s    fdd|j D }t|S )Nc                    r  r5   r   )r8   xr|   r5   r6   r9     r  z-CodeGenerator.visit_Tuple.<locals>.<listcomp>)r	  r^   )rp   rg   ri   r5   r|   r6   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__)rX   r   rw   resub)rp   method_namerN  rO  reverse_method_namer5   r5   r6   _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   rZ   rM  r   formatre   r[  rp   rg   rN  rO  rY  r5   r5   r6   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 }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||
 ||
< qAt| | @ 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 ())rw   r2  r  r   rx   ru   rs   r   rr   rZ   r.  r  get_typerH   keys)rp   rg   rt   
then_block
else_block	then_defs	else_defsrH  r  ir_ret_typesrJ   defs
block_namethen_tyelse_tyr5   r5   r6   visit_then_else_blocks9  sj   





"




z$CodeGenerator.visit_then_else_blocksc                    sX  t | }}|\}}| j }| j }| j| | j|j|| | ||||\ }}}}	}
| j }| j| | rEJ | | j|fdd|D  | j| | rbJ | | j| fdd|D  |
D ]}|	| qrW d    n1 sw   Y  | j
| t|D ]\}}tj|||	| }| || qd S )Nc                       g | ]} | j qS r5   r  r8   nru  r5   r6   r9     r  z4CodeGenerator.visit_if_top_level.<locals>.<listcomp>c                    r}  r5   r  r~  rv  r5   r6   r9     r  )rm   rw   r  r  create_cond_branchr  r|  r3  create_branchadd_argumentr2  rc   r
   corer   rk   r   )rp   condrg   srrt   ip_blockrs  rt  rH  r  rw  endif_blockr2   r   rJ   
new_tensorr5   )rv  ru  r6   visit_if_top_levelp  s0   



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 r5   r&  rw   r7   r|   r5   r6   r9         z.CodeGenerator.visit_if_scf.<locals>.<listcomp>Tr   c                    r}  r5   r  r~  r  r5   r6   r9     r  c                    r}  r5   r  r~  r  r5   r6   r9     r  )rm   r   rw   r  r   r|  r   create_if_opr  merge_block_beforeget_then_blockr  r-  create_yield_opget_else_blockrc   r
   r  r   
get_resultr   )rp   r  rg   r  rt   r   r   last_locrs  rt  rH  r  if_opr   rJ   r  r5   )rv  rp   ru  r6   visit_if_scf  s2   

 
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 )NrU  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 rS   r   r   r5   r5   r6   r         z)CodeGenerator.visit_If.<locals>.<genexpr>)r   testrX   tor
   int1rw   r   r   r   r   r  r  r   rZ   _condition_typesr_  r-   re   r   r   r  )rp   rg   r  contains_returnactive_blockr5   r5   r6   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 )
NrU  zATernary expression with dynamic condition has inconsistent types r  Tr   r  r  c                 s   r  rS   r   r   r5   r5   r6   r     r  z,CodeGenerator.visit_IfExp.<locals>.<genexpr>)$r   r  rX   r  r
   r  rw   rm   r   r  r2  r  r  r   rx   r   r   rZ   r  r&  r  r  r  r  r  r  r  r  r   r  r   r  r   r_  r-   re   )rp   rg   r  r   r  rs  then_valrt  else_valr   ret_type_irr  r5   r5   r6   r     sT   




$#zCodeGenerator.visit_IfExpc                 C      d S rS   r5   r   r5   r5   r6   
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	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   rZ   r   Isr   IsNot_method_name_for_comp_opr   r_  re   r[  )rp   rg   rN  rO  	lhs_value	rhs_valuerY  r5   r5   r6   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.rU  z)' is not (currently) implemented on type )r   operand_method_name_for_unary_opr   rZ   rM  r   re   rX   r   rw   AttributeError)rp   rg   r  rh   r5   r5   r6   visit_UnaryOp  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|k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.)rV   rZ   rX   )rp   rJ   loop_vallive_valr5   r5   r6   _verify_loop_carried_variable$  s    z+CodeGenerator._verify_loop_carried_variablec                    s  t :}|\}} \}}j }j| j| |j j	  j
}|  g }	g }
g }|D ]%}||v r`|| }|| }||| |	| |
|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+|||  qj dd |D  W d    n	1 sCw   Y  t|	D ]\}}tj|!||
| }|j|< |j
|< qL|j"D ]}J d	d S )
Nc                    r  r5   r  r7   r|   r5   r6   r9   M  r  z-CodeGenerator.visit_While.<locals>.<listcomp>c                 S   r  r5   r  r8   rk   r5   r5   r6   r9   N  r  c                    r  r5   r  r7   r|   r5   r6   r9   Q  r  c                    r  r5   )rk   r=   )before_blockr5   r6   r9   Y  r  c                    r  r5   r  r7   r|   r5   r6   r9   \  r  c                 S   r  r5   r  r8   yr5   r5   r6   r9   k  r  FzNot implemented)&rm   r   rw   r  r2  r   r.  r  r   popru   eraser  rZ   r   create_while_opcreate_block_with_parent
get_beforerc   r
   r  r   rk   rr   r   r  r  create_condition_opr  r,  r-  	get_afterr  r  r   r   r  r   )rp   rg   r  rt   ry   r   r  dummy	loop_defsrH  r  	init_argsrJ   r  r  while_opr   r  after_blockyieldsnew_defr  r5   )r  rp   r6   visit_While-  sz   



&

@

zCodeGenerator.visit_Whilec                 C   sJ   |j jjdks	J | |j}| |j}t|r!|j|| jdS || S )NrK  rU  )	r   	__class__re   r   r   slicerX   __getitem__rw   )rp   rg   rN  slicesr5   r5   r6   visit_Subscriptw  s   zCodeGenerator.visit_Subscriptc                    s    fdd|j D S )Nc                    r  r5   r   )r8   dimr|   r5   r6   r9     r  z0CodeGenerator.visit_ExtSlice.<locals>.<listcomp>)dimsr   r5   r|   r6   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rZ||i |}t|j	j
|jj
|jj
}|D ]}t| j|jj<  |j |jD ]	}tj | qMq9d S d }	d }
|tju rz||i |}|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| 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. M}|\}} / \}} j0 } j1|  j23|  |j  j24  |5  g }g }g } j6D ]'}||v r j6| }|| } 7||| |3| |3| |3| q 8||  j9|||dd |D }|	d ur|:d j;|	 |
d ur|:d j;|
  j23|  j1|<d |=  _i  _6t>|D ]\}} ,|tj&-|<d?|d || j@ q |j  j24  g } j6D ]}||v rN|3tj j6|  j q9t|dkrb jAdd |D  |<dB } | C dkstJ d j1|<d |D }|r jE||} jF||} j|jj j)G|  ,|jjtj&-|| W d    n	1 sw   Y  t>|D ]\}} ,|tj&-|H||| j@ q|jD ]}J dd S )Nc                    r  r5   r   r  r|   r5   r6   r9     r  z+CodeGenerator.visit_For.<locals>.<listcomp>c                 3   r   rS   r   r8   keywordr|   r5   r6   r     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  rp  c                 S   r  r5   r  r  r5   r5   r6   r9     r  ztt.num_stagesztt.loop_unroll_factorc                 S   r  r5   r  r  r5   r5   r6   r9     r  z7We use SCF, so the loop body should only have one blockz)Don't know what to do with else after for)Ir   iterr   ri   dictkeywordsr
   static_ranger,  startr   endstepr   rr   r!  r   r  r   r   r   r  r   
num_stagesloop_unroll_factorr-  NumRuntimeErrorrY   r  r  rw   r$   r#   r1   integer_promote_implr&  r'   r  r%   r&   r  create_int_castr%  r   r   rm   r   r  r2  r   r.  r  r  ru   r  r   create_for_opset_attrget_int32_attrget_bodyrs   rc   rk   rZ   r  
get_parentsizeget_induction_var
create_sub
create_addreplace_all_uses_withr  )!rp   rg   IteratorClass	iter_argsiter_kwargsiteratorr  r   r  r  r  lbubr  negative_stepiv_type
iv_ir_typeiv_is_signedivr  rt   ry   r   r  blockr  r  rH  rJ   r  r  for_opfor_op_regionr5   r|   r6   	visit_For  s   


$&&
 "










0


B&
zCodeGenerator.visit_Forc                 C   s0   |  |j}|  |j}|  |j}t|||S rS   )r   lowerupperr  r  )rp   rg   r  r  r  r5   r5   r6   visit_Slice	  s   zCodeGenerator.visit_Slicec                 C   r   rS   )r   r   r   r5   r5   r6   visit_Index  r   zCodeGenerator.visit_Indexc                 C   s   |j | |jfS rS   )rk   r   r   r   r5   r5   r6   visit_keyword  r  zCodeGenerator.visit_keywordc                 C   s:   |  |j}|jd ur|  |jnd}tjj||| jdS )Nr   rU  )r   r  msgr
   r  device_assertrw   )rp   rg   r  r  r5   r5   r6   visit_Assert  s   zCodeGenerator.visit_Assertrh   c                    s  t j|jg R i |  fdd|jD  dd  D  i }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|\}}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|}| j||}| dks|d u rd S | dkrt |!d|S g }t"| D ]}|#t |!|||  qt$|S )Nc                    s   g | ]} | qS r5   r5   )r8   rJ   ri   r5   r6   r9     r:   z2CodeGenerator.call_JitFunction.<locals>.<listcomp>c                 S   s    g | ]}t |r
|nt|qS r5   )rV   r   r  r5   r5   r6   r9          c                 S   s   g | ]
\}}t |r|qS r5   )rY   r8   r   rk   r5   r5   r6   r9      r  c                    s   i | ]}| | qS r5   r5   r=   r  r5   r6   r   !  r  z2CodeGenerator.call_JitFunction.<locals>.<dictcomp>c                    s    g | ]\}}| v rd n|qS rS   r5   r  )
constexprsr5   r6   r9   #  r  c                 S      g | ]	}|d ur|j qS rS   r  r  r5   r5   r6   r9   $      c                 S   r  rS   r  r  r5   r5   r6   r9   %  r  )
r   r   r   r   rb   r   r   r   r   r   r   r   )%inspectgetcallargsrh   rf   rc   rO   re   r   has_functionr
   function_type__globals__r   r   r   r   rb   rw   r   r   r   r   r   	Exceptionr   r   rd   r   r   get_functioncallget_num_resultsr   r  r,  r.  r^   )rp   rh   ri   r   r   r?   arg_vals	arg_typesfn_namer   r   r   r   ro   ecallee_ret_typesymbolcall_opresultsr   r5   )ri   r  r6   call_JitFunction  sN   
zCodeGenerator.call_JitFunctionc           	   
      s*  t  |j} j|}|d ur| |S t fdd|jD } fdd|jD }t|t	r?t
|||  |||S t|drIt|jsOtj|rd ji}t|}d|jv rb |d< z
||i ||W S  ty } z	t jj|d |d }~ww | j v rtt |}||i |S )Nc                 3   r   rS   r   r  r|   r5   r6   r   O  r   z+CodeGenerator.visit_Call.<locals>.<genexpr>c                    r  r5   r   r  r|   r5   r6   r9   P  r  z,CodeGenerator.visit_Call.<locals>.<listcomp>__self__rV  
_generator)r   r   r    statically_implemented_functionsr   r  r  ri   rT   r   rl   r"  hasattrrV   r#  r
   r  
is_builtinrw   r  	signature
parametersr  r   r   rd   r   rI  r.   )	rp   rg   rh   static_implementationkwsri   extra_kwargssigr  r5   r|   r6   r   I  s.   

 


	
zCodeGenerator.visit_Callc                 C   
   t |jS rS   r   r   r   r5   r5   r6   visit_Constanth  rW   zCodeGenerator.visit_Constantrg   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-  rI  r   r   _method_name_for_bool_opr   rZ   rM  r_  re   r[  r`  r5   r5   r6   visit_BoolOpk  s   zCodeGenerator.visit_BoolOplogical_and
logical_orr1  )      c                 C   r.  rS   r/  r   r5   r5   r6   visit_NameConstant{  rW   z CodeGenerator.visit_NameConstantc                 C   r.  rS   )r   r  r   r5   r5   r6   	visit_Num~  rW   zCodeGenerator.visit_Numc                 C   s   t t|S rS   )r   r   literal_evalr   r5   r5   r6   	visit_Str  r`   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   )rw   )	r   r   rX   r   r
   r  permuterw   r   )rp   rg   rN  r5   r5   r6   r     s   zCodeGenerator.visit_Attributec                 C   r  rS   r  r   r5   r5   r6   
visit_Expr  r  zCodeGenerator.visit_Exprc                 C   r  rS   r5   r   r5   r5   r6   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   )r]   rI  rc   rT   r   Constantr(   r   FormattedValue
conversionr   rY   r   rZ   chrr_  AssertionErrorr-   )rp   rg   rI  r   r   conversion_code	evaluatedr5   r5   r6   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   rw   r   r&  r   r   r   rI  rJ  superr   r   r  r   rd   r<   )rp   rg   	last_noder  rN   r  r  r5   r6   r     s0   


$zCodeGenerator.visitc                 C   s   |  |dt|j)Nzunsupported AST node type: {})r   r_  rZ   re   r   r5   r5   r6   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-  ri   r  r1   r   r   rT   r   NotImplementedErrorr  r<   r   r   rd   )rp   rg   	arg_countpassedr   r  r5   r5   r6   execute_static_assert  s*   


z#CodeGenerator.execute_static_assertc                    s   dt jf fdd}|S )Nrg   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 r5   )r   )r8   rJ   r   r5   r5   r6   r     s    z>CodeGenerator.static_executor.<locals>.ret.<locals>.<dictcomp>c                 3   r   rS   r   r  r|   r5   r6   r     r   z=CodeGenerator.static_executor.<locals>.ret.<locals>.<genexpr>c                    s   g | ]	}t  |qS r5   )r   r   r  r|   r5   r6   r9     r  z>CodeGenerator.static_executor.<locals>.ret.<locals>.<listcomp>)r  ri   r   )rp   rg   r+  ri   	python_fnr|   r6   rN     s
   z*CodeGenerator.static_executor.<locals>.ret)r   r   )rY  rN   r5   rX  r6   static_executor  s   zCodeGenerator.static_executorr%  )NFNFNr   )re   r   r   r   r   r   r(   rq   r-  r]   r,  floatintrT   r   r   r   r   updater
   r  device_printminimummaximumr   r   r   r   r   r   r   r   r   r  r   r
  r   r   rC  rD  rF  r   r   r   rQ  rR  rT  r[  ra  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  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__r5   r5   rR  r6   r      s   
 
,2
/	J	
7 5$	J /&(r   c                 C   sJ   d}t | D ]\}}|t|7 }||jv r|d7 }||jv r"|d7 }q|S )Nr   r;   d)rc   r(   
equal_to_1divisibility_16)r(  specializationsuffixr   r   r5   r5   r6   kernel_suffix  s   

r  c                    s  j }fdd  fddj D }j }}	tj }
|	 }|D ]}||
v rB|
| dkrB|| dkrBd||< q,|
 }| }| }|| fdd	j D }t\}}tg |}t|||||	|d|||||d
}|  |j}||_|S )Nc                    s   t | tr j| S | S rS   )rT   r(   rf   index)r   )rh   r5   r6   <lambda>  r  zast_to_ttir.<locals>.<lambda>c                    s   i | ]	\}} ||qS r5   r5   )r8   keyr   )cst_keyr5   r6   r     r  zast_to_ttir.<locals>.<dictcomp>i1r   Tc                    s"   g | ]\}}| j vrt|qS r5   )r?   r   )r8   r   r   )r  r5   r6   r9     r@   zast_to_ttir.<locals>.<listcomp>)r   r?   r   r   r   r   r   r   r   r   r   )attrsr?   r   r  rs   r<   r]   r(  rI  get_constantsfilter_out_constantsget_fn_attrsr]  r   r
   r  r   r   r   r   r   )rh   r  r   r   r   r   r  r?   r   r   tysnew_constantsr   	new_attrsfn_attrsall_constantsr  r   r   r   ro   rN   r5   )r  rh   r  r6   ast_to_ttir  s4   

 
r  )7r   r  rW  r~  rK  r   r   typingr   r   r   r   r   r   r   r   r
   _C.libtritonr   r   r   r   language.corer   r   r   runtime.jitr   r   runtimer   errorsr   r   r   typesr   r!   rO   r   rV   rX   rY   r\   r_   rl   r\  rZ   r  rm   r  r   r   r  r  r5   r5   r5   r6   <module>   sL    $
X        4