
    çg                        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 Z d Z!dede"fdZ#dede"fdZ$dede"fdZ%dede"fdZ&defdZ'd Z(d Z)e"e* e+d          hZ, G d d          Z- G d de j.                  Z/ G d de j.                  Z0d Z1d Z2dS )    N)AnyCallableDictOptionalTupleTypeUnion   )language)ir)	constexprtensor	str_to_ty)_normalize_ty)JITFunction   )CompilationErrorCompileTimeAssertionFailureUnsupportedLanguageConstruct)
ModuleTypec                 8   |                                  rdt          | j                  z   S |                                 r<t          j        j        j        }| j        |k    rdnd}|t          | j
                  z   S |                                 rt          |           S |                                 rIt          | j                  }d                    t          t          | j                            }| d| dS |                                 rd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-   s        Z/var/www/html/ai-engine/env/lib/python3.11/site-packages/triton/compiler/code_generator.pyr   r      s   	yy{{ .Yr}----	yy{{ -*1)V33BO,,,,	~~ 2ww	{{}} !	""S"(++,,      	zz|| s$$$$$    c                 b   d                     d |D                       }d                     fdt                    D                       }|                    dd          }|                    dd          }|                    dd                              d	d          }|  d
| d
| }|S )Nr   c                 ,    g | ]}t          |          S  )r   ).0r/   s     r2   
<listcomp>zmangle_fn.<locals>.<listcomp>'   s    !B!B!BB)B--!B!B!Br3   c                 D    g | ]}| d t          |                    S )c)repr)r7   r   	constantss     r2   r8   zmangle_fn.<locals>.<listcomp>(   s3    !Y!Y!Y!Q"="=il););"="=!Y!Y!Yr3   ._d_'_sq_[]__)r+   sortedreplace)namearg_tysr<   mangled_arg_namesmangled_constantsrets     `   r2   	mangle_fnrK   %   s    !B!B'!B!B!BCC!Y!Y!Y!YviGXGX!Y!Y!YZZ)11#u==)11#v>>)11#s;;CCCMM
=
=&
=
=*;
=
=CJr3   oreturnc                 ,    t          | t                    S N)
isinstancer   rL   s    r2   _is_triton_tensorrR   1   s    a   r3   c                 ,    t          | t                    S rO   )rP   r   rQ   s    r2   _is_constexprrT   5   s    a###r3   c                 t    t          |           o)| j                                         p| j        j        dk    S )Nr   )rR   typer)   numelrQ   s    r2   _is_triton_scalarrX   9   s2    QP):):%:%Oafla>OPr3   c                 :    t          | t          t          f          S rO   )rP   listtuplerQ   s    r2   _is_list_liker\   =   s    a$'''r3   c                 >    t          | t                    r| j        n| S rO   )rP   r   valuerQ   s    r2   _unwrap_if_constexprr_   A   s     I..5177A5r3   c                     |j         rbt          |          D ]T\  }}t          |          s>t          |          s/t	          |j        | d|j         d|j        |          d|           Sd S d S )Nz	Function z= is marked noinline, but was called with non-scalar argument :)noinline	enumeraterT   rX   r   src__name__	arg_names)nodefnargsidxargs        r2   _check_fn_argsrl   E   s    	{ !$ 	 	HC %% .?.D.D 2FD D  D  Djljvwzj{  D  D  B  D  D   	 	r3   c                 P   | }t          |t                    s|j        }t          |t                    |j        j        j        }t          j        |j                  \  }}t          |          D ]3\  }}|                                	                    d          r||z  } n4||fS )Nzdef )
rP   r   rh   __code__co_filenameinspectgetsourcelinesrc   strip
startswith)rh   base_fn	file_namelines
begin_linerj   lines          r2   _get_fn_file_linery   O   s    G+.. * +.. 
#/I.wz::E: u%%  	T::<<""6** 	#JE	 j  r3   c                        e Zd Zd Zd Zd ZdS )enter_sub_regionc                     || _         d S rO   )	generator)selfr}   s     r2   __init__zenter_sub_region.__init__f   s    "r3   c                 N   | j         j                                        | _        | j         j                                        | _        i | j         _        | j         j                                        | _        | j         j        	                                | _
        | j        | j        fS rO   )r}   lscopecopyliveins
local_defs	prev_defsbuilderget_insertion_blockinsert_blockget_insertion_pointinsert_point)r~   s    r2   	__enter__zenter_sub_region.__enter__i   s~    ~,113327799$&! N2FFHH N2FFHH|T...r3   c                     | j         j                            | j                   | j        | j         _        | j        | j         _        d S rO   )r}   r   restore_insertion_pointr   r   r   r   r   )r~   ri   kwargss      r2   __exit__zenter_sub_region.__exit__r   s<    66t7HIII $$(N!!!r3   N)re   
__module____qualname__r   r   r   r6   r3   r2   r{   r{   d   sA        # # #/ / /3 3 3 3 3r3   r{   c                   <   e Zd Zd ZdefdZdefdZdefdZdej	        defdZ
dej        defdZdej        defd	Zdej        defd
Zdej        defdZdej        defdZdej        defdZdej        defdZdej        defdZdej        defdZdS )ContainsReturnCheckerc                     || _         d S rO   )gscope)r~   r   s     r2   r   zContainsReturnChecker.__init__{   s    r3   rM   c                 @    |D ]}|                      |          r dS dS )NTFvisit)r~   bodyss      r2   _visit_stmtsz"ContainsReturnChecker._visit_stmts~   s4     	 	Azz!}} ttur3   c                     t          |t                    rB|j        s;|                                }t	          | j                                      |          S dS NF)rP   r   rb   parser   r   r   )r~   rh   fn_nodes      r2   _visit_functionz%ContainsReturnChecker._visit_function   sO    b+&& 	Er{ 	EhhjjG(55;;GDDDur3   c                 8   d}t          j        |          D ]\  }}t          |t                    r7|D ]3}t          |t           j                  r|p|                     |          }4Qt          |t           j                  r|p|                     |          }|S r   )astiter_fieldsrP   rZ   ASTr   )r~   rg   rJ   r   r^   items         r2   generic_visitz#ContainsReturnChecker.generic_visit   s    -- 	/ 	/HAu%&& /! 6 6D!$00 6!5TZZ%5%56 E37++ /.TZZ..
r3   rg   c                     t          |j        t          j                  rV|j        j        | j        v rA| j        |j        j                 }t          ||j                  }|                     |          S dS | 	                    |j                  S r   )
rP   r^   r   Nameidr   getattrattrr   r   )r~   rg   r^   rh   s       r2   visit_Attributez%ContainsReturnChecker.visit_Attribute   sy     dj#(++ 	z}++DJM2UDI..++B///5zz$*%%%r3   c                     t          |j                  t          j        k    rdS |j        | j        v r'| j        |j                 }|                     |          S dS r   )rV   ctxr   Storer   r   r   )r~   rg   rh   s      r2   
visit_Namez ContainsReturnChecker.visit_Name   sR    >>SY&&57dk!!TW%B''+++ur3   c                     dS )NTr6   r~   rg   s     r2   visit_Returnz"ContainsReturnChecker.visit_Return       tr3   c                     dS r   r6   r   s     r2   visit_Assignz"ContainsReturnChecker.visit_Assign   	     ur3   c                     dS r   r6   r   s     r2   visit_AugAssignz%ContainsReturnChecker.visit_AugAssign   r   r3   c                 6    |                      |j                  S rO   r   r   r   s     r2   visit_Modulez"ContainsReturnChecker.visit_Module         +++r3   c                 6    |                      |j                  S rO   r   r   s     r2   visit_FunctionDefz'ContainsReturnChecker.visit_FunctionDef   r   r3   c                     |                      |j                  }|j        r|p|                      |j                  }|S rO   )r   r   orelse)r~   rg   rJ   s      r2   visit_IfzContainsReturnChecker.visit_If   sB     	**; 	87**4;77C
r3   c                 j    |                      |j                  p|                      |j                  S rO   )r   r   r   r   s     r2   visit_IfExpz!ContainsReturnChecker.visit_IfExp   s)    zz$)$$?

4;(?(??r3   c                 6    |                      |j                  S rO   )r   funcr   s     r2   
visit_Callz ContainsReturnChecker.visit_Call   s    zz$)$$$r3   N)re   r   r   r   boolr   r   r   r   	Attributer   r   r   Returnr   Assignr   	AugAssignr   Moduler   FunctionDefr   Ifr   IfExpr   Callr   r6   r3   r2   r   r   y   s         D    T    	T 	 	 	 	&CM &d & & & &sx D              
CM d    
, , , , , ,,co ,$ , , , ,SV     @	 @d @ @ @ @%sx %D % % % % % %r3   r   c                   p    e Zd ZU 	 	 dededee         dee         fdZ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 Zd Zd Zdedeeef         ddfdZd Zd Z d Z!d Z"d Z#d Z$d Z%d Z&d Z'd Z(d Z)d Z*d  Z+d! Z,d" Z-d# Z.d$ Z/d% Z0e1j2        d&e1j3        d'e1j4        d(e1j5        d)e1j6        d*e1j7        d+e1j8        d,e1j9        d-e1j:        d.e1j;        d/e1j<        d0e1j=        d1iZ>ee?e1j@                 ef         ed2<   d3 ZAd4 ZBd5 ZCd6 ZDd7 ZEd8 ZFd9 ZGe1jH        d:e1jI        d;e1jJ        d<e1jK        d=e1jL        d>e1jM        d?iZNee?e1jO                 ef         ed@<   dA ZPe1jQ        dBe1jR        dCe1jS        dDe1jT        dEiZUee?e1jV                 ef         edF<   dG ZWdH ZXdI ZYdJ ZZdK Z[dL Z\de]eef         fdMZ^defdNZ_dOefdPZ`dQ ZadR ZbdSe1jc        fdTZde1je        dUe1jf        dViZgee?e1jh                 ef         edW<   eijj        dXk     r	dY ZkdZ Zld[ Zmd\ Znd] Zod^ Zpd_ Zq fd`Zrda ZsdSe1jt        ddfdbZudc Zvej        jw        euej        jx         evey          e eve          e eve          iZzee{e|e1jt        gef         f         edd<    xZ}S )fCodeGeneratorNFr   jit_fnfunction_typesru   c                 f   || _         t          j        |          | _        || _        |dz
  | _        | j                            ||d           || j        _        |	| j        _        || j                                        n|| _	        |i n|| _
        || _        || _        t                      | _        || _        || _        || _        || _        || _        d | _        |
|j        n|
| _        || _        g | _        d | _        i | _        |                                 | _        d | _        d| _        d S )Nr   r   F)contextr   r   ru   rw   set_locoptionscodegen_fnscreate_modulemodulefunction_ret_types	prototyper   dictr   
attributesr<   r   function_name	is_kernelcur_nodedebugrb   	scf_stackret_typer   _define_name_lookupdereference_namerh   visiting_arg_default_value)r~   r   r   r   r   r<   r   r   r   r   r   r   r   r   rb   ru   rw   s                    r2   r   zCodeGenerator.__init__   s+    z'**"$q.Y
A666& $/ 6<ndl00222&(6(>""N"ff$"*"&+mW]]
  .06:6N6N6P6P +0'''r3   c                     i | ]
}|j         |S r6   re   r7   r   s     r2   
<dictcomp>zCodeGenerator.<dictcomp>   s    (t(t(t1Q(t(t(tr3   builtin_namespaceprintminmaxc                 8    t          | j        j        ||          S rO   )r   r   rd   )r~   rg   messages      r2   _unsupportedzCodeGenerator._unsupported  s    +DKOT7KKKr3   c                    t                      }| j                            ||          }||u rdS t          |          rdS | j                            di                               |          x}rt	          |          dk    S dS )NFT__annotations__r   )objectr   getrT   r   )r~   rF   absent_markervalas        r2   _is_constexpr_globalz"CodeGenerator._is_constexpr_global  s    koodM22-5 	4 1266::4@@@1 	3 ##{22ur3   c                      dt           f fddt           f fdt                      dt           dt          f fd}|S )NrF   c                 :    j                             | |          S rO   )r   r   )rF   absentr~   s     r2   local_lookupz7CodeGenerator._define_name_lookup.<locals>.local_lookup  s    ;??4000r3   c                 B   j                             | |          }||u s| j        v st          |          t          k    st          |t                    st          |dd          s~t          |dd                              d          sZt          |t          j
                  s@                    |           s+j        s$t          j                            dd          dk    r|S t          t!          j        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   rV   r   rP   r   r   rs   r   r"   r   r   osenviron	NameErrortextwrapdedentrE   )rF   r   r   r~   s      r2   global_lookupz8CodeGenerator._define_name_lookup.<locals>.global_lookup  s9   +//$//C vt555CyyJ..!#{33 /s$8%@@ / sL"55@@ARSS / "#x~66	 /
 0066 / 6 / z~~&JCPPTWWW
 !4/3!4 !4 !4 5 5 6=WT35G5GI I Ir3   rM   c                 x    }j         j        fD ]} || |          }||ur|c S t          |  d          )Nz is not defined)r   r   r  )rF   r   lookup_functionr^   r   r  r   r~   s       r2   name_lookupz6CodeGenerator._define_name_lookup.<locals>.name_lookup8  se    "F#/@V@Z#Z ! !'f55&& LLL 't444555r3   )r&   r   r   )r~   r  r   r  r   s   ` @@@r2   r   z!CodeGenerator._define_name_lookup  s    	1s 	1 	1 	1 	1 	1 	1	I 	I 	I 	I 	I 	I 	I8 	6c 	6c 	6 	6 	6 	6 	6 	6 	6 	6 	6 r3   rF   r^   rM   c                 .    || j         |<   || j        |<   dS )z This function:
            called by visit_Assign() & visit_FunctionDef() to store left value (lvalue)
        1. record local defined name (FIXME: should consider control flow)
        2. store tensor in self.lvalue
        N)r   r   )r~   rF   r^   s      r2   	set_valuezCodeGenerator.set_valueB  s"     "D %r3   c                 n    | j                                         }| j                                         }||fS rO   )r   get_locr   )r~   locips      r2   _get_insertion_point_and_locz*CodeGenerator._get_insertion_point_and_locK  s3     l""$$\--//3wr3   c                 n    | j                             |           | j                             |           d S rO   )r   r   r   )r~   r  r  s      r2   _set_insertion_point_and_locz*CodeGenerator._set_insertion_point_and_locS  s4    ,,R000S!!!!!r3   c                     t          |          s|g}|D ]4}|                     |           t          |t          j                  r d S 5d S rO   )r\   r   rP   r   r   )r~   stmtsstmts      r2   visit_compound_statementz&CodeGenerator.visit_compound_statementZ  sf    U## 	GE 	 	DJJt $
++ 	 	r3   c                 F    t           j                            | |           d S rO   r   NodeVisitorr   r   s     r2   r   zCodeGenerator.visit_Modulef       %%dD11111r3   c                 j                           |j                  }|J  fd|j        D             }|S )Nc                 :    g | ]}                     |          S r6   r   )r7   r1   r~   s     r2   r8   z,CodeGenerator.visit_List.<locals>.<listcomp>l  #    555C

3555r3   )r   r   elts)r~   rg   r   r&  s   `   r2   
visit_ListzCodeGenerator.visit_Listi  s?    jj""{{{555549555r3   c                 V                          |j                  }|' j                            g            t          j        }nt          |t                    rN fd|D             }d |D             } j                            d |D                        t          |          }nLt          j        	                    | j                  } j                            |j
        g           |j        } j        	| _        d S  j        |k    rt          d j         d|           d S )Nc                 Z    g | ]'}t           j                            |j                  (S r6   )r   core
_to_tensorr   )r7   vr~   s     r2   r8   z.CodeGenerator.visit_Return.<locals>.<listcomp>z  s-    WWW(-221dlCCWWWr3   c                     g | ]	}|j         
S r6   rV   r7   r,  s     r2   r8   z.CodeGenerator.visit_Return.<locals>.<listcomp>{  s    444A444r3   c                     g | ]	}|j         
S r6   handler/  s     r2   r8   z.CodeGenerator.visit_Return.<locals>.<listcomp>|  s    ;;;1ah;;;r3   zInconsistent return types:  and )r   r^   r   rJ   r   voidrP   r[   r*  r+  r2  rV   r   	TypeError)r~   rg   	ret_valueret_ty
ret_values	ret_typesrJ   s   `      r2   r   zCodeGenerator.visit_Returnp  s4   JJtz**	
 LR   ]FF	5)) 	WWWWYWWWJ44444IL;;
;;;<<<9%%FF-**9dlCCCLcj\***XF = "DMMM]f$$V$-VVfVVWWW %$r3   c                    |                      |j                  \  }}| j        r|                     |d          t	          |j        j                  D ]\  }}|j        j        | dz
           }|j        }|j        }t          j	        |t          j
                              }	|t          j        |	g|          }
nt          j        |	||          }
	 | j        rJ d| _        |                      |
           d| _        # d| _        w xY w| j        rdnd	}| j                            | j        | j        | j                            | j                  || j                  | _        | j                            | j                   | j                                        }g }d
}t	          |          D ]\  }}|| j        v rL| j        |         }t3          |          st5          | j        |                   }|                    |           Z|| j        v r/| j        |         D ]!\  }}| j                            |||           "|                    t=          | j                            |          | j        j        |                              |dz  }| j                                         }tC          ||          D ]\  }}| "                    ||           | j        #                    |           | $                    |j%                   | j&        | j&        tN          j(        k    r,tN          j(        | _&        | j        )                    g            ntU          | j&        tV                    rVtY          | j&                  | j        _-        | j        .                    | j                            | j                             nI| j&        g| j        _-        | j        .                    | j                            | j                             |r| j        /                    |           | j        0                                 d S )Nz,nested function definition is not supported.r   r   r   targetsr^   )targetr^   
annotationTFpublicprivater   )1r   ri   rh   r   rc   defaultsr?  rk   r   r   r   r   	AnnAssignr   r   r   get_or_insert_functionr   r   r   to_irrb   	push_backadd_entry_blockr<   rT   r   appendr   set_arg_attrr   param_typesr   zipr  set_insertion_point_to_startr  r   r   r   r4  rJ   rP   r[   rZ   r9  
reset_typeset_insertion_point_to_endfinalize)r~   rg   rf   kwarg_namesr   default_valuearg_noder?  rF   	st_target	init_node
visibilityentry
arg_valuesrj   arg_namecstr^   	insert_pt	arg_values                       r2   r   zCodeGenerator.visit_FunctionDef  s
   !%DI!6!6	;7 	Z##D*XYYY )$)*< = = 	8 	8A}y~qb1f-H!,J<DDcikk:::I!J	{-PPP		M-\fggg	8::::26/

9%%%27//%/7777 "&>XXY
,55dk4CU6:n6J6J4<6X6XZdfjfsu udg&&&''))
$Y// 	 	KAxDN""nQ'$S)) 7#DN1$566C!!#&&&'''+q'9 ? ?e,,S$>>>>!!&c):):DN<VWZ<["\"\]]]qL4466	#&y*#=#= 	0 	0HiNN8Y////11%888%%di000= DMX]$B$B$MDMLR     $-// G+/+>+>(""4>#7#7#E#EFFFF,0M?(""4>#7#7#E#EFFF 	?L33I>>>s   %D	Dc                     g }|j         D ]}||                     |          gz  }|                     |j                  }||fS rO   )ri   r   kwarg)r~   rg   rf   rk   rP  s        r2   visit_argumentszCodeGenerator.visit_arguments  sO    	9 	+ 	+C$**S//**IIjj,,+%%r3   c                 P    t           j                            | |           |j        S rO   )r   r!  r   rk   r   s     r2   	visit_argzCodeGenerator.visit_arg  s!    %%dD111xr3   c                 ~   |                      |j                  }|                      |j                  }|                      |j                  }|t          k    rP|| 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   
ValueErrorrT   r   )r~   rg   r?  r>  r^   s        r2   visit_AnnAssignzCodeGenerator.visit_AnnAssign  s    ZZ00
DK((

4:&&""$$ F "E "E "E F F F '' )!%(("'DK;v&&  &&&r3   c                 J   g }|j         D ]}||                     |          gz  }t          |          dk    r|                     |d          |d         }|                     |j                  }t          |          s|g}t          |          s|g}t          j        f}t          ||          D ]p\  }}t          |          }|Dt          |          s5t          ||          s%t          j                            || j                  }|                     ||           qd S )Nr   z2simultaneous multiple assignment is not supported.r   )r=  r   lenr   r^   r\   r   r"   rK  r_   rR   rP   r*  r+  r   r  )	r~   rg   _namesr>  namesvaluesnative_nontensor_typesrF   r^   s	            r2   r   zCodeGenerator.visit_Assign  s=   l 	+ 	+Ftzz&))**FFv;;??##D*^___q	DJ''U## 	GEV$$ 	XF"*.!3uf-- 	( 	(KD%(//E $U++ !e%;<< ! !00EENN4''''	( 	(r3   c                 4   |j         j        }t          j        |t          j                              }t          j        ||j        |j                  }t          j        |j         g|          }| 	                    |           | 
                    |          S )Nr;  r<  )r>  r   r   r   LoadBinOpopr^   r   r   r   )r~   rg   rF   lhsrhsassigns         r2   r   zCodeGenerator.visit_AugAssign  sz    {~h$CHJJ///iTWdj11T[M===

6$$T***r3   c                     t          |j                  t          j        k    r|j        S |                     |j                  S rO   )rV   r   r   r   r   r   r   s     r2   r   zCodeGenerator.visit_Name  s4    >>SY&&7N$$TW---r3   c                 F    t           j                            | |           d S rO   r   r   s     r2   visit_StorezCodeGenerator.visit_Store  r"  r3   c                 F    t           j                            | |           d S rO   r   r   s     r2   
visit_LoadzCodeGenerator.visit_Load  r"  r3   c                 H      fd|j         D             }t          |          S )Nc                 :    g | ]}                     |          S r6   r   )r7   xr~   s     r2   r8   z-CodeGenerator.visit_Tuple.<locals>.<listcomp>  s#    111!

1111r3   )r&  r[   )r~   rg   ri   s   `  r2   visit_TuplezCodeGenerator.visit_Tuple  s*    1111ty111T{{r3   c                    t          |          r  t          ||          || j                  S t          |          r6t          j        dd|          } t          ||          || j                  S  t          ||          |          S )N_builderz__(.*)__z__r\1__)rR   r   r   resub)r~   method_namern  ro  reverse_method_names        r2   _apply_binary_methodz"CodeGenerator._apply_binary_method  s    S!! 	I,73,,S4<HHHHS!! 	Q"$&j+"N"N473 344S4<PPPP(wsK((---r3   c                 Z   |                      |j                  }|                      |j                  }| j                            t          |j                            }|3|                     |d                    |j        j	                            | 
                    |||          S )Nz8AST binary operator '{}' is not (currently) implemented.)r   leftright_method_name_for_bin_opr   rV   rm  r   formatre   r  r~   rg   rn  ro  r  s        r2   visit_BinOpzCodeGenerator.visit_BinOp  s    jj##jj$$266tDG}}EE##D$^$e$efjfmfv$w$wy y y((c3???r3   __add____sub____mul____truediv____floordiv____mod____pow__
__lshift__
__rshift____and____or____xor__r  c                    | j                             |           |                     |j                   | j                                         }| j                                        }i }|j        r| j                             |           |                                | _        i | _        |                     |j                   | j                                        }| j                                         }g }g }g }	|D ]&}
|df|dffD ]Q\  }}|
|v rH||
         j	        ||
         j	        k    s,J d|
 d||
         j	         d| d||
         j	                     R|
|v s|
|v r|
                    |
           |
                    |
|v r||
         j	        n||
         j	                   |	
                    |
|v r||
         j                                        n||
         j                                                   |
|v r|
|vr||
         ||
<   |
|v r|
|vr||
         ||
<   (|                                |                                z  D ]}
|
|v r||
         j	        }||
         j	        }||k    sJ d|
 d| d	| d
            |
                    |
           |
                    |           |	
                    ||
         j                                                   |||||||	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   rL  r  r   r   r   r   r   r   rV   rH  r2  get_typekeys)r~   rg   r   
then_block
else_block	then_defs	else_defsrg  r9  ir_ret_typesrF   defs
block_namethen_tyelse_tys                  r2   visit_then_else_blocksz$CodeGenerator.visit_then_else_blocks7  s   11*===%%di000\5577
O((**		; 	<L55jAAA!,,..DK DO))$+666,,..I99;;J 	 	0 	0D&/%89f:M$N X X j4<<:?gdm.@@@@Xd X XAS X X#-X XFJ4joX X A@@ y  DI$5$5T"""  9J9J4!5!5PYZ^P_Pdeee##$-J. J.IdO$:$C$C$E$E$E3<T?3I3R3R3T3TV V V y  T%:%:")$-	$y  T%:%:")$-	$ NN$$y~~'7'77 
	C 
	CDu}}o*Go*Gg%%%.t . .' . .#*. . . &%% LLW%%%	$ 6 ? ? A ABBBB)ZUI|[[r3   c                    d}t          |           5 }|\  }}| j                                        }| j                                        }| j                                        }	| j                            |           | j                            |j        ||           |                     ||||          \  }}}
}}| j                            |           |                                r*|                                rd}|	                                 |	                                s)|r'| j        
                    |	fd|
D                        | j                            |           |	                                s)|r'| j        
                    |	fd|
D                        |r|D ]}|	                    |           d d d            n# 1 swxY w Y   |r~| j                            |	           t          |
          D ]V\  }}t          j                            |	                    |          ||                   }|                     ||           Ud S d S )NTFc                 *    g | ]}|         j         S r6   r1  r7   nr  s     r2   r8   z4CodeGenerator.visit_if_top_level.<locals>.<listcomp>  !    8\8\8\QR19L8\8\8\r3   c                 *    g | ]}|         j         S r6   r1  r7   r  r  s     r2   r8   z4CodeGenerator.visit_if_top_level.<locals>.<listcomp>  r  r3   )r{   r   create_blockrN  create_cond_branchr2  r  
has_returnerasehas_terminatorcreate_branchadd_argumentrL  rc   r   r*  r   rk   r  )r~   condrg   has_endif_blocksrr   ip_blockr  r  endif_blockrg  r9  r  r/   r   rF   
new_tensorr  r  s                    @@r2   visit_if_top_levelz CodeGenerator.visit_if_top_leveln  s   d## 	1r "GX2244J2244J,3355KL33H===L++DKZPPP ++D':zRR YIy*j%L L33J???$$&& $:+@+@+B+B $"'!!###,,.. ^? ^**;8\8\8\8\V[8\8\8\]]]L33J???,,.. ^? ^**;8\8\8\8\V[8\8\8\]]] 1& 1 1B,,R00003	1 	1 	1 	1 	1 	1 	1 	1 	1 	1 	1 	1 	1 	1 	14  	1L55kBBB$U++ 1 14%]11+//!2D2DiPQlSS
tZ0000	1 	11 1s   F4GGGc                     t                     5 }|\  }}                                 \  }} j                                        }|j        r j                                        nd }	                     ||||	          \  }}	}
}}                     ||            j                             fd|D             |j        d          }|	                    |
                                            j                            |
                                           t          |
          dk    r& j                            fd|
D                        |j        s|                                }	n'|		                    |                                            j                            |                                           t          |
          dk    r& j                            fd|
D                        d d d            n# 1 swxY w Y   t          |
          D ]T\  }}t           j                            |                    |          ||                   }                     ||           Ud S )Nc                 D    g | ]}|                     j                  S r6   rE  r   r7   r/   r~   s     r2   r8   z.CodeGenerator.visit_if_scf.<locals>.<listcomp>  s'    .Z.Z.Z"rxx/E/E.Z.Z.Zr3   Tr   c                 *    g | ]}|         j         S r6   r1  r  s     r2   r8   z.CodeGenerator.visit_if_scf.<locals>.<listcomp>       -Q-Q-Qail.A-Q-Q-Qr3   c                 *    g | ]}|         j         S r6   r1  r  s     r2   r8   z.CodeGenerator.visit_if_scf.<locals>.<listcomp>  r  r3   )r{   r  r   r  r   r  r  create_if_opr2  merge_block_beforeget_then_blockrN  re  create_yield_opget_else_blockrc   r   r*  r   
get_resultr  )r~   r  rg   r  r   r   r  last_locr  r  rg  r9  if_opr   rF   r  r  r  s   `               @@r2   visit_if_scfzCodeGenerator.visit_if_scf  s   d## 	SrJGQ<<>>LB2244J8<M22444J++D':zRR NIy*j%A --b(;;;L--.Z.Z.Z.ZPY.Z.Z.Z\`\gimnnE))%*>*>*@*@AAAL33E4H4H4J4JKKK5zzA~~,,-Q-Q-Q-Q5-Q-Q-QRRR; F"1133

--e.B.B.D.DEEEL33E4H4H4J4JKKK5zzA~~,,-Q-Q-Q-Q5-Q-Q-QRRR)	S 	S 	S 	S 	S 	S 	S 	S 	S 	S 	S 	S 	S 	S 	S, !'' 	- 	-GAt!--e.>.>q.A.A9Q<PPJNN4,,,,	- 	-s   G/HHHc           	         |                      |j                  }t          |          r|                    t          j        | j                  }t          | j                                       |          }| j	        r|r| 
                    |d          | j	        s|s|                     ||           d S |                     ||           d S t          |          }t          |          t          vr^| 
                    |d                    d                    d t          D                       t          |          j                            |r|                     |j                   d S |                     |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              3   $   K   | ]}|j         V  d S rO   r   r   s     r2   	<genexpr>z)CodeGenerator.visit_If.<locals>.<genexpr>  $      !G!G!*!G!G!G!G!G!Gr3   )r   testrR   tor   int1r   r   r   r   r   r  r  r_   rV   _condition_typesr  r+   re   r  r   r   )r~   rg   r  contains_returns       r2   r   zCodeGenerator.visit_If  s   zz$)$$T"" 	;778=4<7@@D3DK@@FFtLLO~ 4/ 4'' PQ Q Q  4 4!!$-----''d33333'--DDzz!111''krr		!G!G6F!G!G!GGGT

+- -. . .  ;--di88888--dk:::::r3   c           	      |   |                      |j                  }t          |          rZ|                    t          j        | j                  }t          |           5  |                                 \  }}| j        	                                }| j        
                    |           t          j                            |                      |j                  | j                  }| j                                        }| j        	                                }| j        
                    |           t          j                            |                      |j                  | j                  }| j                                        }|                     ||           |j        |j        k    sJ d|j         d|j                     |j        }	|	t          j        k    r|	                    | j                  gng }
| j                            |
|j        d          }|                    |                                           |
rL| j                            |                                           | j                            |j        g           | j                            |                                           |                    |                                           |
rL| j                            |                                           | j                            |j        g           |
r3t          j                            |                    d          |	          nd cd d d            S # 1 swxY w Y   d S t;          |          }t#          |          t<          vr^|                     |d                     d!                    d t<          D                       t#          |          j"                            |r|                      |j                  S |                      |j                  S )	Nr{  zAternary expression with dynamic condition has inconsistent types r3  Tr   r  r  c              3   $   K   | ]}|j         V  d S rO   r   r   s     r2   r  z,CodeGenerator.visit_IfExp.<locals>.<genexpr>  r  r3   )#r   r  rR   r  r   r  r   r{   r  r  rL  r*  r+  r   r   r   r  rV   r4  rE  r  r2  r  r  rN  r  r  r   r  r_   r  r   r  r+   re   )r~   rg   r  r  r  r  then_valr  else_valr   ret_type_irr  s               r2   r   zCodeGenerator.visit_IfExp  s   zz$)$$T"" 1	/778=4<7@@D!$'' !d !d#@@BBH!\6688
99*EEE#=33DJJty4I4I4<XX!\==??
!\6688
99*EEE $=33DJJt{4K4KT\ZZ!\==??
11"h???}555{X`Xe{{ltly{{ 655#=@HHM@Y@Yx~~dl;;<<_a11+t{DQQ--e.B.B.D.DEEE DL;;E<P<P<R<RSSSL00(/1BCCC778L8L8N8NOOO--e.B.B.D.DEEE DL;;E<P<P<R<RSSSL00(/1BCCCNYcx}++E,<,<Q,?,?JJJ_cC!d !d !d !d !d !d !d !d !d !d !d !d !d !d !d !d !d !dF (--D Dzz!111''krr		!G!G6F!G!G!GGGT

+- -. . .  /zz$),,,zz$+...s    LM88M<?M<c                     d S rO   r6   r   s     r2   
visit_PasszCodeGenerator.visit_Pass  s    r3   c                 *   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    rt          ||u          S | j                            t          |j        d                             }|9|                     |d                    |j        d         j                            |                     |||          S )Nr   z1simultaneous multiple comparison is not supportedr   z<AST comparison operator '{}' is not (currently) implemented.)re  comparatorsopsr   r   r  r_   rV   r   Isr   IsNot_method_name_for_comp_opr   r  re   r  )r~   rg   rn  ro  	lhs_value	rhs_valuer  s          r2   visit_ComparezCodeGenerator.visit_Compare  sa   D$%%**s48}}/A/A##D*]^^^jj##jj)!,--(--	(--	&&Y)3444	))Yi7888377TXa[8I8IJJ##T[[\`\def\g\pqqs s s((c3???r3   __eq____ne____lt____le____gt____ge__r  c           
         |                      |j                  }| j                            t	          |j                            }|$|                     |d|j        j         d          t          |          r t          ||          | j
                  S 	  t          ||                      S # t          $ r/ |                     |d| dt	          |          j                   w xY w)NzAST unary operator 'z!' is not (currently) implemented.r{  z)' is not (currently) implemented on type )r   operand_method_name_for_unary_opr   rV   rm  r   re   rR   r   r   AttributeError)r~   rg   r  rh   s       r2   visit_UnaryOpzCodeGenerator.visit_UnaryOp  s	   **T\**+//TW>>:##D*tAQ*t*t*tuuuW%% 	?'77B''>>>>	t'77B''))) 	t 	t 	t##rRrrZ^_fZgZgZprrt t t	ts   B4 49C-__neg____pos____not__
__invert__r  c           
      <
    t                     5 }|\  }}                                 \  }} j                                        } j                            |            j                            |                                |j                    j        	                                  j
        }|                                 g }	g }
g }|D ]}||v rt          ||                   sJ d| d            t          ||                   sJ d| d            ||         j        ||         j        k    s*J d| d||         j         d||         j         d            |	                    |           |
                    ||         j                   |                    ||                                         ||            j                             fd|
D             d	 |D                       } j                            |                                 fd
|
D                        j                                       t%          |	          D ][\  }}t&          j                                                |          |
|                    j        |<    j        |          j
        |<   \                     |j                  } j                                        j                            |j        fdt;          t=          |                    D                         j                            |                                 fd|
D                       } j                            |           t%          |	          D ][\  }}t&          j                            |                    |          |
|                    j        |<    j        |          j
        |<   \ j                            |                                |j                    j        	                                  j
        }g }|D ]!}||v r|                    ||                    " j                             d |D                        d d d            n# 1 swxY w Y   t%          |	          D ]R\  }}t&          j                            |!                    |          |
|                   }| j        |<   | j
        |<   S|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                 D    g | ]}|                     j                  S r6   r  r  s     r2   r8   z-CodeGenerator.visit_While.<locals>.<listcomp>J  s(    4`4`4`PRRXXdl5K5K4`4`4`r3   c                     g | ]	}|j         
S r6   r1  r7   rk   s     r2   r8   z-CodeGenerator.visit_While.<locals>.<listcomp>K  s    4U4U4UCSZ4U4U4Ur3   c                 D    g | ]}|                     j                  S r6   r  r  s     r2   r8   z-CodeGenerator.visit_While.<locals>.<listcomp>N  s(    AmAmAm]_"((4<BXBXAmAmAmr3   c                 :    g | ]}                     |          S r6   )rk   )r7   r   before_blocks     r2   r8   z-CodeGenerator.visit_While.<locals>.<listcomp>V  s(    :n:n:nST<;K;KA;N;N:n:n:nr3   c                 D    g | ]}|                     j                  S r6   r  r  s     r2   r8   z-CodeGenerator.visit_While.<locals>.<listcomp>Y  s(    @l@l@l\^$,AWAW@l@l@lr3   c                     g | ]	}|j         
S r6   r1  r7   ys     r2   r8   z-CodeGenerator.visit_While.<locals>.<listcomp>h  s    )C)C)Cq!()C)C)Cr3   FzNot implemented)&r{   r  r   r  rL  r   rH  r  r   popr   r  rR   rV   r  create_while_opcreate_block_with_parent
get_beforerc   r   r*  r   rk   r   r   r  rN  create_condition_opr2  rangere  	get_afterr  r  r   r   r!  r   )r~   rg   r  r   r   r  r  dummy	loop_defsrg  r9  	init_argsrF   while_opr   r  after_blockyieldsnew_defr  r  s   `                   @r2   visit_WhilezCodeGenerator.visit_While'  s   d## @	Er$&!G\<<>>LB L--//EL55e<<<N!!$'''))$)444N   IKKMMM EII! 4 47??,Yt_==mm?m[_?m?m?mmmm,WT];;kk=kY]=k=k=kkkk$T?/74=3EEEEL L LQUI[ L L1:41EL L L FEE LL&&&$$Yt_%9:::$$WT]333--b(;;;|334`4`4`4`V_4`4`4`4U4U94U4U4UW WH  <@@ATATAVAVAmAmAmAmclAmAmAmo oLL55lCCC$U++ : :4$,M$8$89I9I!9L9LiXYl$[$[D!(,D(9%%::di((DL33LAAAL,,T[:n:n:n:nX]^abk^l^lXmXm:n:n:nooo,??@R@R@T@T@l@l@l@lbk@l@l@ln nK L55kBBB$U++ : :4$,M$8$89K9KYWX\$Z$ZD!(,D(9%%N!!$'''))$)444N   IF! 3 37??MM)D/222L(()C)CF)C)C)CDDDA@	E @	E @	E @	E @	E @	E @	E @	E @	E @	E @	E @	E @	E @	E @	EF !'' 	, 	,GAtm**8+>+>q+A+A9Q<PPG 'DK$+DOD!!K 	6 	6D+++++	6 	6s   Q>RR #R c                     |j         j        j        dk    sJ |                     |j                  }|                     |j                  }t          |          r|                    || j                  S ||         S )Nrk  r{  )	r   	__class__re   r   r^   slicerR   __getitem__r   )r~   rg   rn  slicess       r2   visit_SubscriptzCodeGenerator.visit_Subscriptt  st    x!*f4444jj$$DJ''S!! 	B??6DL?AAA6{r3   c                 *      fd|j         D             S )Nc                 :    g | ]}                     |          S r6   r   )r7   dimr~   s     r2   r8   z0CodeGenerator.visit_ExtSlice.<locals>.<listcomp>}  r%  r3   )dimsr   s   ` r2   visit_ExtSlicezCodeGenerator.visit_ExtSlice|  s    5555495555r3   c                                           |j        j                  } fd|j        j        D             }t	           fd|j        j        D                       }|t          j        k    r ||i |}t          |j	        j
        |j        j
        |j        j
                  }|D ]g}t          |           j        |j        j        <                        |j                   |j        D ]"}t(          j                             |           #hd S d }	|t          j        u r% ||i |}|j	        }
|j        }|j        }|j        }	n|t          u rt1          |          dk    r|d         n&                      t)          j        d                    }
t1          |          dk    r|d         n$                      |j        j        d                   }t1          |          dk    r|d         n&                      t)          j        d                    }nt5          d          d}t7          |          r&|j
        dk     rt          |j
                   }d}||
}}
t          j                            |
 j                  }
t          j                            | j                  }t          j                            | j                  }|
j                                         r2|j                                         r|j                                         s(tC          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[                     5 }|\  }} .                                \  }} j        /                                } j        0                    |            j1        2                    |                                |j                    j1        3                                 |4                                 g }g }g } j5        D ],}||v r$tm           j5        |                   sJ | d            tm          ||                   sJ  j5        |         j7        ||         j7        k    s/J d| d||         j7         d j5        |         j7         d            |2                    |           |2                    t          j                            ||          j                             |2                    t          j                             j5        |          j                             . 8                    ||            j        9                    |
||d |D                       }|	.|:                    d j        ;                    |	                      j1        2                    |            j        0                    |<                    d                     |=                                 _        i  _5        t}          |          D ]m\  }} +                    |t          j        ,                    |<                    d          ?                    |dz             ||         j7                             n                     |j                    j1        3                                 g } j5        D ]I}||v rC|2                    t          j                             j5        |          j                             Jt1          |          dk    r$ j        @                    d |D                        |<                    d          A                                }|B                                dk    s
J d             j        0                    |<                    d                     |C                                }|r6 j        D                    ||          } j        E                    ||
          } j        |j        j                 j(        F                    |            +                    |j        j        t          j        ,                    ||                     d d d            n# 1 swxY w Y   t}          |          D ]W\  }} +                    |t          j        ,                    |G                    |          ||         j7                             X|j        D ]}J d            d S )Nc                 :    g | ]}                     |          S r6   r   r7   rk   r~   s     r2   r8   z+CodeGenerator.visit_For.<locals>.<listcomp>  s#    ???TZZ__???r3   c              3   B   K   | ]}                     |          V  d S rO   r   r7   keywordr~   s     r2   r  z*CodeGenerator.visit_For.<locals>.<genexpr>  s/      QQ74::g..QQQQQQr3   r   r   r
   zAOnly `range` and `static_range` iterators are currently supportedFTz0For loop bounds and step must all be ints, are (r  r  z is not tensorr  r  r  r  c                     g | ]	}|j         
S r6   r1  r  s     r2   r8   z+CodeGenerator.visit_For.<locals>.<listcomp>  s    >_>_>_csz>_>_>_r3   ztt.num_stagesc                     g | ]	}|j         
S r6   r1  r  s     r2   r8   z+CodeGenerator.visit_For.<locals>.<listcomp>  s    -G-G-G1ah-G-G-Gr3   z7We use SCF, so the loop body should only have one blockz)Don't know what to do with else after for)Hr   iterr   ri   r   keywordsr   static_ranger  startr^   endstepr   r   r>  r   r  r   r   r   r!  r   
num_stagesre  NumRuntimeErrorrT   r*  r+  r   r"   r!   r5  semanticinteger_promote_implrE  r%   r#   r$   r2  create_int_castcreate_undefr  r   r{   r  r  rL  r   rH  r  r  r   rR   rV   r  create_for_opset_attrget_int32_attrget_bodyr   rc   rk   r  
get_parentsizeget_induction_var
create_sub
create_addreplace_all_uses_withr  )r~   rg   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  rg  rF   for_opfor_op_regions   `                             r2   	visit_ForzCodeGenerator.visit_For  s	   

49>22????	???	QQQQdi>PQQQQQH111$}i?;??H !5x|7I8=K^__L! > >.7llDKN+--di888 K > >DO11$====>F
HN**$}i?;??H BB=D!,JJe## "%Y!!3!31CGAJJ9O9OB!$Y!!3!31DINSTDU9V9VB#&y>>A#5#59Q<<4::cgajj;Q;QDDbccc 	4:>>dj[))D MB]%%b$,77]%%b$,77}''dl;;x   	v(9(9 	vARARATAT 	vtrxtt[][cttgkgqtttuuu#8828LL#88$*MM]]4<00
-1D1O1VVYY{\))"j,GG\))"j,GG|++D*lKK\&&z22t{~x}';';B'H'HIIId## @	Nr$&!G\<<>>LB L--//EL55e<<<N!!$'''))$)444N   KKMMM IFE a a7??,T_T-BCC\\E\E\E\\\\,WT];;;;;?4059KKKKL L LQUI[ L L151F1KL L L LKK
 LL&&&$$X]%=%=gdmT\%Z%Z[[[MM(-":":4?4;PRVR^"_"_``` --b(;;;\//B>_>_U^>_>_>_``F%1L1LZ1X1XYYYN!!$'''L55fooa6H6HIII!,,..DK DO$U++ j j4tX]%9%9&//!:L:L:P:PQRUVQV:W:WY_`aYbYg%h%hiiii))$)444N   F a a7??MM(-":":4?4;PRVR^"_"_``` 6{{Q,,-G-G-G-G-GHHH"OOA..99;;M %%''1,,,.g,,, L55fooa6H6HIII))++B 5\,,R44\,,R44K'.DDRHHHNN4;>8=+?+?G+L+LMMMA@	N @	N @	N @	N @	N @	N @	N @	N @	N @	N @	N @	N @	N @	N @	NF !'' 	] 	]GAtNN4!5!5f6G6G6J6JFSTIN![![\\\\K 	6 	6DEEEEE	6 	6s   T)gggc                     |                      |j                  }|                      |j                  }|                      |j                  }t	          |||          S rO   )r   lowerupperr*  r  )r~   rg   rL  rM  r*  s        r2   visit_SlicezCodeGenerator.visit_Slice  sM    

4:&&

4:&&zz$)$$UE4(((r3   c                 6    |                      |j                  S rO   )r   r^   r   s     r2   visit_IndexzCodeGenerator.visit_Index  s    zz$*%%%r3   c                 D    |j         |                     |j                  fS rO   )rk   r   r^   r   s     r2   visit_keywordzCodeGenerator.visit_keyword  s    xDJ////r3   c                     | j         sd S |                     |j                  }|j        |                     |j                  nd}t          j                            ||| j                  S )Nr  r{  )r   r   r  msgr   r*  device_assertr   )r~   rg   r  rT  s       r2   visit_AssertzCodeGenerator.visit_Assert  sb    z 	Fzz$)$$&*h&:djj"""}**4t|*LLLr3   rh   c                    t          j        |j        gR i |fd|j        D             d D             t	                      }d t                    D             fdD             }fdt                    D             d D             }d D             }t          |j        ||          }| j        	                    |          st          j        g |          }	|j        }
t          |          \  }}|j        | j        n|j        }t          | j        |	|
||| j        ||| j        |j        ||| j        j        | j        j        |          }	 |                    |                                           n3# t0          $ r&}t3          | j        j        | j        d           |d }~ww xY w|j        }|| j        |<   n| j        |         }| j                            |          }| j                            ||          }|                                 d	k    s|d S |                                 d
k    r#tC          |"                    d	          |          S g }tG          |                                           D ]>}|$                    tC          |"                    |          ||                              ?tK          |          S )Nc                      g | ]
}|         S r6   r6   )r7   rF   ri   s     r2   r8   z2CodeGenerator.call_JitFunction.<locals>.<listcomp>  s    444tT
444r3   c                 N    g | ]"}t          |          r|nt          |          #S r6   )rR   r   r  s     r2   r8   z2CodeGenerator.call_JitFunction.<locals>.<listcomp>  s0    RRRc(--A9S>>RRRr3   c                 6    g | ]\  }}t          |          |S r6   )rT   )r7   r   rk   s      r2   r8   z2CodeGenerator.call_JitFunction.<locals>.<listcomp>  s)    LLLFAss9K9KLaLLLr3   c                 "    i | ]}||         S r6   r6   )r7   r   ri   s     r2   r   z2CodeGenerator.call_JitFunction.<locals>.<dictcomp>   s    444AQQ444r3   c                 &    g | ]\  }}|v rd n|S rO   r6   )r7   r   rk   
constexprss      r2   r8   z2CodeGenerator.call_JitFunction.<locals>.<listcomp>"  s)    OOOVQZSOOOr3   c                      g | ]}||j         S rO   r1  r  s     r2   r8   z2CodeGenerator.call_JitFunction.<locals>.<listcomp>#  s    BBB3#/CJ///r3   c                      g | ]}||j         S rO   r.  r  s     r2   r8   z2CodeGenerator.call_JitFunction.<locals>.<listcomp>$  s    AAA#SXr3   )
r   r   r   r   rb   ru   rw   r   r   r   r   r   )&rp   getcallargsrh   rf   r   rc   rK   re   r   has_functionr   function_type__globals__ry   r   r   r   r   rb   r   r   r   r   r   	Exceptionr   r   rd   r   r   get_functioncallget_num_resultsr   r  r  rH  r[   )r~   rh   ri   r   r   r<   arg_vals	arg_typesfn_namer   r   ru   rw   r   r}   ecallee_ret_typesymbolcall_opresultsr   r]  s     `                  @r2   call_JitFunctionzCodeGenerator.call_JitFunction  s   "25:4:::6::4444r|444RRTRRRVV
LLiooLLL
4444444	OOOOyOOOBB$BBBAAAAA	BKI>>{''00 	? .r9==I^F$5b$9$9!Iz"$("2DJJE%dlIvzS\eiep-/wW[Wn/1{i\f.2l.BPTP\Phpuw w wIT

++++ T T T&t{tLLRSST (0O/>D#G,,"5g>O))'22,##FH55""$$))_-D4$$&&!++',,Q//AAA G7224455 R Rvg&8&8&;&;_Q=OPPQQQQ>>!s   'F 
F6!F11F6c                     t                               |j                            } j                            |          }| | |          S t           fd|j        D                       } fd|j        D             }|t          j	        j
        u r	 j        sd S t          |t                    r(t          |||                                |||          S t!          |d          rt#          |j                  st          j	                            |          rpt           j                  }t+          j        |          }d|j        v r |d<   	  ||i ||S # t0          $ r!}t3           j        j        |d           |d }~ww xY w| j                                        v rt=          t           |          } ||i |S )Nc              3   B   K   | ]}                     |          V  d S rO   r   r!  s     r2   r  z+CodeGenerator.visit_Call.<locals>.<genexpr>N  s/      DD74::g&&DDDDDDr3   c                 :    g | ]}                     |          S r6   r   r  s     r2   r8   z,CodeGenerator.visit_Call.<locals>.<listcomp>O  r%  r3   __self__r{  
_generator)r_   r   r    statically_implemented_functionsr   r   r&  ri   r   r*  rU  r   rP   r   rl   rp  hasattrrR   rt  
is_builtinr   rp   	signature
parametersrd  r   r   rd   r   rh  r,   )	r~   rg   rh   static_implementationkwsri   extra_kwargssigrk  s	   `        r2   r   zCodeGenerator.visit_CallH  s   !$**TY"7"788 $ E I I" M M ,((t444DDDDdmDDDDD555549555,,,: b+&& 	84T***((T3777B
## 	K(9"+(F(F 	K8=KcKcdfKgKg 	K666L#B''Cs~---1\*	Kr47<73777 K K K 't{dCCJK '..0000+T22Dr43s   	E" "
F,FFc                 *    t          |j                  S rO   r   r^   r   s     r2   visit_ConstantzCodeGenerator.visit_Constantj  s    $$$r3   rg   c                    t          |j                  dk    r|                     |d          |                     |j        d                   }|                     |j        d                   }| j                            t          |j                            }|3|                     |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.)re  rh  r   r   _method_name_for_bool_opr   rV   rm  r  re   r  r  s        r2   visit_BoolOpzCodeGenerator.visit_BoolOpm  s    t{q  ##vx x xjjQ((jjQ((377TWFF##QXXY]Y`Yijjl l l((c3???r3   logical_and
logical_orr  )      c                 *    t          |j                  S rO   r  r   s     r2   visit_NameConstantz CodeGenerator.visit_NameConstant}  s    TZ(((r3   c                 *    t          |j                  S rO   )r   r  r   s     r2   	visit_NumzCodeGenerator.visit_Num  s    TV$$$r3   c                 D    t          t          j        |                    S rO   )r   r   literal_evalr   s     r2   	visit_StrzCodeGenerator.visit_Str  s    S-d33444r3   c                     |                      |j                  }t          |          r2|j        dk    r't          j                            |d| j                  S t          ||j                  S )NT)r   r   )r   )	r   r^   rR   r   r   r.  permuter   r   )r~   rg   rn  s      r2   r   zCodeGenerator.visit_Attribute  se    jj$$S!! 	TyC(00fdl0SSSsDI&&&r3   c                 F    t           j                            | |           d S rO   r   r   s     r2   
visit_ExprzCodeGenerator.visit_Expr  r"  r3   c                     d S rO   r6   r   s     r2   visit_NoneTypezCodeGenerator.visit_NoneType  r   r3   c           
         t          |j                  }t          |          D ]\  }}t          |t          j                  rt          |j                  ||<   8t          |t          j                  r|j	        }| 
                    |j                  }t          |          s3|                     |dt          t          |                    z             |dk     rdndt          |          z   dz                       |j                  ||<   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  )rZ   rh  rc   rP   r   Constantr&   r^   FormattedValue
conversionr   rT   r   rV   chrr  AssertionErrorr+   )r~   rg   rh  r   r^   conversion_code	evaluateds          r2   visit_JoinedStrzCodeGenerator.visit_JoinedStr  sD   dk""!&)) 	w 	wHAu%.. w,,q		E3#566 
w"'"2 JJu{33	$Y// 0++xd9oo../0 0 0 &5q%8%8TTdSEYEY>Y\_>_gghqhwxxq		$%a%h%himnsitit%u%uvvvwwvr3   c           	         |d S t          j                    5  t          j        dt                     t          j        dt                     | j        }| j                                        }|| _        t          |d          r\t          |d          rL| j        	                    | j
        | j        |j        z   |j                   | j                                        }	 t                                          |          }nJ# t           $ r  t"          $ r3}t!          | j        j        | j        t)          |                    d d }~ww xY w|r!|| _        | j        	                    |           |cd d d            S # 1 swxY w Y   d S )Nignorelineno
col_offset)warningscatch_warningssimplefilterDeprecationWarningPendingDeprecationWarningr   r   r  rw  r   ru   rw   r  r  superr   r   rd  r   rd   r;   )r~   rg   	last_noder  rJ   rk  r  s         r2   r   zCodeGenerator.visit  s   <F$&& 	 	 !(,>???!(,EFFFI|++--H DMtX&& 274+F+F 2$$T^T_t{5RTXTcddd<//11ZggmmD))#    Z Z Z 't{tAwwOOUYYZ  / )$$X...1	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	s6   CE?"!DE?E.EE'E??FFc                 x    |                      |d                    t          |          j                            )Nzunsupported AST node type: {})r   r  rV   re   r   s     r2   r   zCodeGenerator.generic_visit  s1    &E&L&LTRVZZM`&a&abbbr3   c                 8   t          |j                  }d|cxk     rdk    rn nt          |j                  rt          d          t	          |                     |j        d                             }t          |t                    st          d          |s|dk    rd}nN	 |                     |j        d                   }n,# t          $ r}dt          |          z   dz   }Y d }~nd }~ww xY 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: >)re  ri   r&  r5  r_   r   rP   r   NotImplementedErrorrd  r;   r   r   rd   )r~   rg   	arg_countpassedr   rk  s         r2   execute_static_assertz#CodeGenerator.execute_static_assert  s<   	NN	I"""""""""s4='9'9"[\\\%djj1&>&>??&$'' 	% D    		dA~~X"jj166GG  X X XG$q''QTWWGGGGGGX .dkotEYZaEbEbcccts   % C 
C/C**C/c                 ,     dt           j        f fd}|S )Nrg   c                      d  fd|j         D             D             } fd|j        D             }t           |i |          S )Nc                 4    i | ]\  }}|t          |          S r6   )r_   )r7   rF   r^   s      r2   r   z>CodeGenerator.static_executor.<locals>.ret.<locals>.<dictcomp>  s7       D% *511  r3   c              3   B   K   | ]}                     |          V  d S rO   r   r!  s     r2   r  z=CodeGenerator.static_executor.<locals>.ret.<locals>.<genexpr>  s/      #U#UGDJJw$7$7#U#U#U#U#U#Ur3   c                 T    g | ]$}t                              |                    %S r6   )r_   r   r  s     r2   r8   z>CodeGenerator.static_executor.<locals>.ret.<locals>.<listcomp>  s,    OOOc(C99OOOr3   )r&  ri   r   )r~   rg   r|  ri   	python_fns   `   r2   rJ   z*CodeGenerator.static_executor.<locals>.ret  ss     #U#U#U#Ut}#U#U#U  C POOOTYOOODYY444555r3   )r   r   )r  rJ   s   ` r2   static_executorzCodeGenerator.static_executor  s1    	6CH 	6 	6 	6 	6 	6 	6 
r3   rv  )NNFNFNr   )~re   r   r   r   r   r   r&   r   re  rZ   r  floatintrP   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`  rc  r   r   r   rs  ru  ry  r  r  r   AddSubMultDivFloorDivModPowLShiftRShiftBitAndBitOrBitXorr  r   operatorr  r  r  r   r   r  r  EqNotEqLtLtEGtGtEr  cmpopr  USubUAddNotInvertr  unaryopr  r  r  rJ  rN  rP  r   rR  rV  rp  r   r  BoolOpr  AndOrr  boolopsysversion_infor  r  r  r   r  r  r  r   r   r   r  r  static_assertstatic_printr   rv  r   r   __classcell__)r  s   @r2   r   r      s|         jnMN$0 $0al $0X`aeXf$0,4SM$0 $0 $0 $0L )u(tdESXZ]_ikr@s(t(t(ttCH~ttt	(-,-	 !	 !   L L L  , , ,\&c &%	0A*B &t & & & &  " " "
 
 
2 2 2  X X X4A A AF& & &  ' ' '"( ( (,+ + +. . .
2 2 22 2 2  . . .@ @ @ 	)n
L
L
I	8
I>T$s|"4c"9:   5\ 5\ 5\n"1 "1 "1J- - -6; ; ;63/ 3/ 3/j  @ @ @$ 	#)Xsvx(TWTZ\dfifmow<d4	?C#78   t t t 	)SXy#'9cjR^?tD$5s$:;   K6 K6 K6Z  6 6 6D6 D6 D6L) ) )& & &0U38_ 0 0 0 0MC M M M M-"; -" -" -" -"^        D% % %
@ 
@ 
@ 
@ 
@ >AWmUXU[]i<jd4
#3S#89jjj
&  	) 	) 	)	% 	% 	%	5 	5 	5' ' '2 2 2    $    :c c c#( t    ,
 
 
 	#%:"OOE$:$:__S!!__S!!	Q$d68SXJO3L+L&M       r3   r   c                     d}t          |           D ]3\  }}|t          |          z  }||j        v r|dz  }||j        v r|dz  }4|S )Nr  r:   d)rc   r&   
equal_to_1divisible_by_16)ry  specializationsuffixr   r   s        r2   kernel_suffixr    sh     F)$$  1#a&&)))cMF...cMFMr3   c                     j         } fdfdj                                        D             } j                                        }                               }t          j                                                  fd|j	        D             }	d |j
        D             }
|                                }|                    |	           fdj                                        D             }t                     \  }}t          j        g |          }t          ||||| |
d||||          }|                                                                |j        }||_        |S )Nc                 f    t          | t                    rj                            |           n| S rO   )rP   r&   rf   index)r   rh   s    r2   <lambda>zast_to_ttir.<locals>.<lambda>  s+    As1C1CJ**1--- r3   c                 .    i | ]\  }} |          |S r6   r6   )r7   keyr^   cst_keys      r2   r   zast_to_ttir.<locals>.<dictcomp>   s'    XXXeuXXXr3   c                 :    i | ]}||v r|         d k    rdndS )i1Tr   r6   )r7   ktyss     r2   r   zast_to_ttir.<locals>.<dictcomp>  s3    ]]]qQSSVt^^]]]r3   c                     i | ]}|d gS ))ztt.divisibility   r6   )r7   r  s     r2   r   zast_to_ttir.<locals>.<dictcomp>  s    MMM!,-MMMr3   c                 F    g | ]\  }}|j         vt          |          S r6   )r<   r   )r7   r  r,  r  s      r2   r8   zast_to_ttir.<locals>.<listcomp>
  s1    ppp$!QaWeWoNoNo1NoNoNor3   T)
r   r<   r   r   r   r   ru   rw   r   r   )attrsr<   itemsrc  r   r;   rZ   ry  rh  r  r  r  ry   r   rb  r   r   r   r   r   )rh   r  r   r   r   r  r<   r   r   new_constants	new_attrsall_constantsri  ru   rw   r   r}   rJ   r  r  s   ``                @@r2   ast_to_ttirr    s    EJJJJGXXXX~7O7U7U7W7WXXXI^  ""FGGN++M
~'..00
1
1C]]]]EL\]]]MMMu7LMMMINN$$M'''pppp.*B*H*H*J*JpppI-b11Iz&r955Igy=hu%'IYb)3WR]_ _ _I OOBHHJJ

CCKJr3   )3r   rp   r}  r  r  r	  r  typingr   r   r   r   r   r   r	   r  r   _C.libtritonr   r   r   r   runtime.jitr   runtimer   errorsr   r   r   typesr   r   rK   r   rR   rT   rX   r\   r_   rl   ry   r  rV   r  r{   r!  r   r   r  r  r6   r3   r2   <module>r     s   



  				 



  				  D D D D D D D D D D D D D D D D D D             3 3 3 3 3 3 3 3 3 3 ' ' ' ' ' ' ! ! ! ! ! ! a a a a a a a a a a      % % %$	 	 	! ! ! ! ! !$S $T $ $ $ $Q Q Q Q Q Q(S (T ( ( ( (6C 6 6 6 6  ! ! !$ #ttDzz* 3 3 3 3 3 3 3 3*X% X% X% X% X%CO X% X% X%vX X X X XCO X X Xv 
 
 
    r3   