
    sg                        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 Z2y)    N)AnyCallableDictOptionalTupleTypeUnion   )language)ir)	constexprtensor	str_to_ty)_normalize_ty)JITFunction   )CompilationErrorCompileTimeAssertionFailureUnsupportedLanguageConstruct)
ModuleTypec                 "   | j                         rdt        | j                        z   S | j                         rOt        j
                  j                  j                  }| j                  |k(  rdnd}|t        | j                        z   S | j                         rt        |       S | j                         rFt        | j                        }dj                  t        t        | j                               }| d| dS | j#                         ryJ 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)tyr$   prefixeltr-   s        Q/var/www/html/venv/lib/python3.12/site-packages/triton/compiler/code_generator.pyr   r      s    	yy{Yr}}---	yy{**11))V3BOO,,,	~~2w	{{}		"S"((+,awa  	zz|$$$5    c                 |   dj                  |D cg c]  }t        |       c}      }dj                  t        |      D cg c]  }| dt        ||           c}      }|j	                  dd      }|j	                  dd      }|j	                  dd      j	                  dd      }|  d	| d	| }|S c c}w c c}w )
Nr   c._d_'_sq_[]__)r+   r   sortedreprreplace)namearg_tys	constantsr/   mangled_arg_namesr   mangled_constantsrets           r2   	mangle_fnrF   %   s    '!BB)B-!BCviGX!Y!QCqil);(<"=!YZ)11#u=)11#v>)11#s;CCCMF"&'r*;)<
=CJ "C!Ys   B4B9oreturnc                 "    t        | t              S N)
isinstancer   rG   s    r2   _is_triton_tensorrM   1   s    a  r3   c                 "    t        | t              S rJ   )rK   r   rL   s    r2   _is_constexprrO   5   s    a##r3   c                     t        |       xr6 | j                  j                          xs | j                  j                  dk(  S )Nr   )rM   typer)   numelrL   s    r2   _is_triton_scalarrS   9   s2    QP):%:%Oafflla>OPr3   c                 .    t        | t        t        f      S rJ   )rK   listtuplerL   s    r2   _is_list_likerW   =   s    a$''r3   c                 >    t        | t              r| j                  S | S rJ   )rK   r   valuerL   s    r2   _unwrap_if_constexprrZ   A   s     I.1775A5r3   c                     |j                   rbt        |      D ]S  \  }}t        |      rt        |      rt	        |j
                  | d|j                   d|j                  |    d|        y y )Nz	Function z= is marked noinline, but was called with non-scalar argument :)noinline	enumeraterO   rS   r   src__name__	arg_names)nodefnargsidxargs        r2   _check_fn_argsrg   E   s    	{{!$ 	HC %.?.D2FFD},ijljvjvwzj{i||}  B  ~C  D 	 r3   c                 d   | }t        |t              s|j                  }t        |t              s|j                  j                  j                  }t        j                  |j                        \  }}t        |      D ].  \  }}|j                         j                  d      s&||z  } ||fS  ||fS )Nzdef )
rK   r   rc   __code__co_filenameinspectgetsourcelinesr^   strip
startswith)rc   base_fn	file_namelines
begin_linere   lines          r2   _get_fn_file_linert   O   s    G+.** +.

##//I..wzz:E: u% 	T::<""6*#Jj  	 j  r3   c                       e Zd Zd Zd Zd Zy)enter_sub_regionc                     || _         y rJ   )	generator)selfrx   s     r2   __init__zenter_sub_region.__init__f   s	    "r3   c                    | j                   j                  j                         | _        | j                   j                  j                         | _        i | j                   _        | j                   j                  j                         | _        | j                   j                  j                         | _
        | j                  | j                  fS rJ   )rx   lscopecopyliveins
local_defs	prev_defsbuilderget_insertion_blockinsert_blockget_insertion_pointinsert_point)ry   s    r2   	__enter__zenter_sub_region.__enter__i   s    ~~,,11322779$&! NN22FFH NN22FFH||T....r3   c                     | j                   j                  j                  | j                         | j                  | j                   _        | j                  | j                   _        y rJ   )rx   r   restore_insertion_pointr   r~   r|   r   r   )ry   rd   kwargss      r2   __exit__zenter_sub_region.__exit__r   s@    66t7H7HI $$(NN!r3   N)r`   
__module____qualname__rz   r   r    r3   r2   rv   rv   d   s    #/3r3   rv   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j2                  defdZdej6                  defdZy)ContainsReturnCheckerc                     || _         y rJ   )gscope)ry   r   s     r2   rz   zContainsReturnChecker.__init__{   s	    r3   rH   c                 8    |D ]  }| j                  |      s y yNTFvisit)ry   bodyss      r2   _visit_stmtsz"ContainsReturnChecker._visit_stmts~   s$     	Azz!}	 r3   c                     t        |t              r@|j                  s4|j                         }t	        | j
                        j                  |      S yNF)rK   r   r]   parser   r   r   )ry   rc   fn_nodes      r2   _visit_functionz%ContainsReturnChecker._visit_function   s:    b+&r{{hhjG(5;;GDDr3   c                 4   d}t        j                  |      D ]}  \  }}t        |t              r8|D ]2  }t        |t         j                        s|xs | j                  |      }4 Nt        |t         j                        si|xs | j                  |      } |S r   )astiter_fieldsrK   rU   ASTr   )ry   rb   rE   r   rY   items         r2   generic_visitz#ContainsReturnChecker.generic_visit   s    - 	/HAu%&! 6D!$0!5TZZ%56 E377+.TZZ.	/ 
r3   rb   c                 Z   t        |j                  t        j                        rm|j                  j                  | j
                  v rJ| j
                  |j                  j                     }t        ||j                        }| j                  |      S y| j                  |j                        S r   )
rK   rY   r   Nameidr   getattrattrr   r   )ry   rb   rY   rc   s       r2   visit_Attributez%ContainsReturnChecker.visit_Attribute   sw     djj#((+zz}}+DJJMM2UDII.++B//zz$**%%r3   c                     t        |j                        t        j                  k(  ry|j                  | j
                  v r*| j
                  |j                     }| j                  |      S yr   )rQ   ctxr   Storer   r   r   )ry   rb   rc   s      r2   
visit_Namez ContainsReturnChecker.visit_Name   sO    >SYY&77dkk!TWW%B''++r3   c                      y)NTr   ry   rb   s     r2   visit_Returnz"ContainsReturnChecker.visit_Return       r3   c                      yr   r   r   s     r2   visit_Assignz"ContainsReturnChecker.visit_Assign        r3   c                      yr   r   r   s     r2   visit_AugAssignz%ContainsReturnChecker.visit_AugAssign   r   r3   c                 8    | j                  |j                        S rJ   r   r   r   s     r2   visit_Modulez"ContainsReturnChecker.visit_Module         ++r3   c                 8    | j                  |j                        S rJ   r   r   s     r2   visit_FunctionDefz'ContainsReturnChecker.visit_FunctionDef   r   r3   c                     | j                  |j                        }|j                  r|xs | j                  |j                        }|S rJ   )r   r   orelse)ry   rb   rE   s      r2   visit_IfzContainsReturnChecker.visit_If   s=     		*;;7**4;;7C
r3   c                 r    | j                  |j                        xs | j                  |j                        S rJ   )r   r   r   r   s     r2   visit_IfExpz!ContainsReturnChecker.visit_IfExp   s'    zz$))$?

4;;(??r3   c                 8    | j                  |j                        S rJ   )r   funcr   s     r2   
visit_Callz ContainsReturnChecker.visit_Call   s    zz$))$$r3   N)r`   r   r   rz   boolr   r   r   r   	Attributer   r   r   Returnr   Assignr   	AugAssignr   Moduler   FunctionDefr   Ifr   IfExpr   Callr   r   r3   r2   r   r   y   s   D T 	T 	&CMM &d &sxx D     
CMM d 
, , ,,coo ,$ ,SVV  @		 @d @%sxx %D %r3   r   c                      e Zd ZU 	 	 dbdedee   dee   fdZee	e
eeeefD  ci c]  }|j                   | c}} Zeeef   ed<   ej%                  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" Z0e1jd                  d#e1jf                  d$e1jh                  d%e1jj                  d&e1jl                  d'e1jn                  d(e1jp                  d)e1jr                  d*e1jt                  d+e1jv                  d,e1jx                  d-e1jz                  d.iZ>ee?e1j                     ef   ed/<   d0 ZAd1 ZBd2 ZCd3 ZDd4 ZEd5 ZFd6 ZGe1j                  d7e1j                  d8e1j                  d9e1j                  d:e1j                  d;e1j                  d<iZNee?e1j                     ef   ed=<   d> ZPe1j                  d?e1j                  d@e1j                  dAe1j                  dBiZUee?e1j                     ef   edC<   dD ZWdE ZXdF ZYdG ZZdH Z[dI Z\de]eef   fdJZ^defdKZ_dLefdMZ`dN ZadO ZbdPe1j                  fdQZde1j                  dRe1j                  dSiZgee?e1j                     ef   edT<   eij                  dUk  r	dV ZkdW ZldX ZmdY ZndZ Zod[ Zpd\ Zqfd]Zrd^ ZsdPe1j                  ddfd_Zud` Zvej(                  j                  euej(                  j                   evey      e eve      e eve      iZzee{e|e1j                  gef   f   eda<   xZ}S c c}} w )cCodeGeneratorNjit_fnfunction_typesrp   c                 x   || _         t        j                  |      | _        || _        |dz
  | _        | j                  j                  ||d       || j                  _        |	| j                  _        || j                  j                         n|| _	        |i n|| _
        || _        || _        t               | _        || _        || _        || _        || _        || _        d | _        |
|j*                  n|
| _        || _        g | _        d | _        i | _        | j5                         | _        d | _        d| _        y )Nr   r   F)contextr   r   rp   rr   set_locoptionscodegen_fnscreate_modulemodulefunction_ret_types	prototyper   dictr|   
attributesrB   r   function_name	is_kernelcur_nodedebugr]   	scf_stackret_typer   _define_name_lookupdereference_namerc   visiting_arg_default_value)ry   r   r   r   r   rB   r   r   r   r   r   r   r   r   r]   rp   rr   s                    r2   rz   zCodeGenerator.__init__   s    zz'*"$q.Y
A6& $/ 6<ndll002&(6(>"N"f$"*"&+mW]]
  .06:6N6N6P +0'r3   builtin_namespaceprintminmaxc                 D    t        | j                  j                  ||      S rJ   )r   r   r_   )ry   rb   messages      r2   _unsupportedzCodeGenerator._unsupported  s    +DKKOOT7KKr3   c                     t               }| j                  j                  ||      }||u ryt        |      ry| j                  j                  di       j                  |      x}rt	        |      dk(  S y)NFT__annotations__r   )objectr   getrO   r   )ry   r@   absent_markervalas        r2   _is_constexpr_globalz"CodeGenerator._is_constexpr_global  sl    kkoodM2- 126::4@@1@ #{22r3   c                 |     dt         f fddt         f fdt               dt         dt        f fd}|S )Nr@   c                 <    j                   j                  | |      S rJ   )r|   r   )r@   absentry   s     r2   local_lookupz7CodeGenerator._define_name_lookup.<locals>.local_lookup  s    ;;??400r3   c                    j                   j                  | |      }||u s| j                  v st        |      t        k(  st        |t              st        |dd      svt        |dd      j                  d      sZt        |t        j                        s@j                  |       s/j                  s#t        j                  j                  dd      dk(  r|S t        t!        j"                  d	|  d
      j%                  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   rQ   r   rK   r   r   rn   r   r"   r   r   osenviron	NameErrortextwrapdedentr?   )r@   r   r   ry   s      r2   global_lookupz8CodeGenerator._define_name_lookup.<locals>.global_lookup  s    ++//$/C vt555CyJ.!#{3s$8%@sL"5@@ARS!#x~~6006 66zz~~&JCPTWW
 %//3f 51!4 5 6=WT35GI Ir3   rH   c                     }j                   j                  fD ]  } || |      }||us|c S  t        |  d      )Nz is not defined)r   r   r  )r@   r   lookup_functionrY   r   r  r   ry   s       r2   name_lookupz6CodeGenerator._define_name_lookup.<locals>.name_lookup8  sW    "F#/@V@V@Z@Z#Z !'f5& L! tfO455r3   )r&   r   r   )ry   r
  r   r  r   s   ` @@@r2   r   z!CodeGenerator._define_name_lookup  sF    	1s 	1	I 	I8 	6c 	6c 	6 	6 r3   r@   rY   rH   c                 @    || j                   |<   || j                  |<   y)z This function:
            called by visit_Assign() & visit_FunctionDef() to store left value (lvalue)
        1. record local defined name (FIXME: should consider control flow)
        2. store tensor in self.lvalue
        N)r|   r   )ry   r@   rY   s      r2   	set_valuezCodeGenerator.set_valueB  s      "D %r3   c                 r    | j                   j                         }| j                   j                         }||fS rJ   )r   get_locr   )ry   locips      r2   _get_insertion_point_and_locz*CodeGenerator._get_insertion_point_and_locK  s1     ll""$\\--/3wr3   c                 p    | j                   j                  |       | j                   j                  |       y rJ   )r   r   r   )ry   r  r  s      r2   _set_insertion_point_and_locz*CodeGenerator._set_insertion_point_and_locS  s&    ,,R0S!r3   c                     t        |      s|g}|D ]/  }| j                  |       t        |t        j                        s/ y  y rJ   )rW   r   rK   r   r   )ry   stmtsstmts      r2   visit_compound_statementz&CodeGenerator.visit_compound_statementZ  s?    U#GE 	DJJt $

+	r3   c                 D    t         j                  j                  | |       y rJ   r   NodeVisitorr   r   s     r2   r   zCodeGenerator.visit_Modulef      %%dD1r3   c                     | j                  |j                        }|J |j                  D cg c]  }| j                  |       }}|S c c}w rJ   )r   r   elts)ry   rb   r   r1   r  s        r2   
visit_ListzCodeGenerator.visit_Listi  sG    jj"{{+/995C

355 6s   A
c                 &   | j                  |j                        }|,| j                  j                  g        t        j
                  }nt        |t              r|D cg c],  }t        j                  j                  || j                        . }}|D cg c]  }|j                   }}| j                  j                  |D cg c]  }|j                   c}       t        |      }n\t        j                  j                  || j                        }| j                  j                  |j                  g       |j                  }| j                  || _        y | j                  |k7  rt        d| j                   d|       y c c}w c c}w c c}w )NzInconsistent return types:  and )r   rY   r   rE   r   voidrK   rV   core
_to_tensorrQ   handler   	TypeError)ry   rb   	ret_valueret_tyv
ret_values	ret_typesrE   s           r2   r   zCodeGenerator.visit_Returnp  s3   JJtzz*	
 LLR ]]F	5)MVW(--221dllCWJW)34A4I4LL
;1ahh;<9%F--**9dllCCLLcjj\*XXF == "DM]]f$9$--fXVWW % X4;s   1FF	Fc                 z	   | j                  |j                        \  }}| j                  r| j                  |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| _        | j                  |
       d| _         | j                  rdnd	}| j                  j!                  | j"                  | j$                  | j&                  j)                  | j                        || j*                        | _        | j"                  j-                  | j                         | j                  j/                         }g }d
}t	        |      D ]  \  }}|| j0                  v rD| j0                  |   }t3        |      st5        | j0                  |         }|j7                  |       X|| j8                  v r4| j8                  |   D ]"  \  }}| j                  j;                  |||       $ |j7                  t=        | j                  j                  |      | j&                  j>                  |                |dz  } | j                  jA                         }tC        ||      D ]  \  }}| jE                  ||        | j                  jG                  |       | jI                  |jJ                         | jL                  | jL                  tN        jP                  k(  r1tN        jP                  | _&        | j                  jS                  g        ntU        | jL                  tV              rctY        | jL                        | j&                  _-        | j                  j]                  | j&                  j)                  | j                               nZ| jL                  g| j&                  _-        | j                  j]                  | j&                  j)                  | j                               |r| j                  j_                  |       | j                  ja                          y # d| _        w xY w)Nz,nested function definition is not supported.r   r   r   targetsrY   )targetrY   
annotationTFpublicprivater   )1r   rd   rc   r   r^   defaultsr0  rf   r   r   r   r   	AnnAssignr   r   r   get_or_insert_functionr   r   r   to_irr]   	push_backadd_entry_blockrB   rO   r   appendr   set_arg_attrr   param_typesr   zipr  set_insertion_point_to_startr  r   r   r   r!  rE   rK   rV   rU   r*  
reset_typeset_insertion_point_to_endfinalize)ry   rb   ra   kwarg_namesr   default_valuearg_noder0  r@   	st_target	init_node
visibilityentry
arg_valuesre   arg_namecstrY   	insert_pt	arg_values                       r2   r   zCodeGenerator.visit_FunctionDef  s   !%DII!6	;77##D*XYY )$))*<*< = 	8A}yy~~qb1f-H!,,J<<DDciik:I!JJ	{-P	MM-\fg	8::::26/

9%27/	8$ "&XY
,,55dkk4CUCU6:nn6J6J4<<6XZdfjfsfsudgg&'')
$Y/ 	KAxDNN"nnQ'$S)#DNN1$56C!!#&''+q'9 ?e,,S$>?!!&c):DNN<V<VWZ<["\]q	 LL446	#&y*#= 	0HiNN8Y/	011%8%%dii0== DMMX]]$B$MMDMLLR  $--/+/+>(""4>>#7#7#EF,0MM?(""4>>#7#7#EFLL33I>[ 38/s   4&R11	R:c                     g }|j                   D ]  }|| j                  |      gz  } | j                  |j                        }||fS rJ   )rd   r   kwarg)ry   rb   ra   rf   rA  s        r2   visit_argumentszCodeGenerator.visit_arguments  sL    	99 	+C$**S/**I	+jj,+%%r3   c                 Z    t         j                  j                  | |       |j                  S rJ   )r   r  r   rf   r   s     r2   	visit_argzCodeGenerator.visit_arg  s    %%dD1xxr3   c                 x   | j                  |j                        }| j                  |j                        }| j                  |j                        }|t        k(  rP|| j
                  v rt        | d      t        |      st	        |      }|| j
                  |<   | j
                  |   S | j                  |      S )Nz4 is already defined. constexpr cannot be reassigned.)	r   r0  r/  rY   r   r|   
ValueErrorrO   r   )ry   rb   r0  r/  rY   s        r2   visit_AnnAssignzCodeGenerator.visit_AnnAssign  s    ZZ0
DKK(

4::&"$ F8 ,D "E F F '!%("'DKK;;v&&  &&r3   c                    g }|j                   D ]  }|| j                  |      gz  } t        |      dkD  r| j                  |d      |d   }| j                  |j                        }t        |      s|g}t        |      s|g}t        j                  f}t        ||      D ]e  \  }}t        |      }|At        |      s6t        ||      s*t        j                  j                  || j                        }| j                  ||       g y )Nr   z2simultaneous multiple assignment is not supported.r   )r.  r   lenr   rY   rW   r   r"   r<  rZ   rM   rK   r"  r#  r   r  )	ry   rb   _namesr/  namesvaluesnative_nontensor_typesr@   rY   s	            r2   r   zCodeGenerator.visit_Assign  s    ll 	+Ftzz&)**F	+v;?##D*^__q	DJJ'U#GEV$XF"*..!3uf- 	(KD%(/E $U+e%;< 00ENN4'	(r3   c                 ^   |j                   j                  }t        j                  |t        j                               }t        j
                  ||j                  |j                        }t        j                  |j                   g|      }| j                  |       | j                  |      S )Nr,  r-  )r/  r   r   r   LoadBinOpoprY   r   r   r   )ry   rb   r@   lhsrhsassigns         r2   r   zCodeGenerator.visit_AugAssign  sr    {{~~hh$CHHJ/iiTWWdjj1T[[M=

6$$T**r3   c                     t        |j                        t        j                  k(  r|j                  S | j                  |j                        S rJ   )rQ   r   r   r   r   r   r   s     r2   r   zCodeGenerator.visit_Name  s4    >SYY&77N$$TWW--r3   c                 D    t         j                  j                  | |       y rJ   r  r   s     r2   visit_StorezCodeGenerator.visit_Store  r  r3   c                 D    t         j                  j                  | |       y rJ   r  r   s     r2   
visit_LoadzCodeGenerator.visit_Load  r  r3   c                 r    |j                   D cg c]  }| j                  |       }}t        |      S c c}w rJ   )r  r   rV   )ry   rb   xrd   s       r2   visit_TuplezCodeGenerator.visit_Tuple  s0    '+yy1!

111T{ 2s   4c                     t        |      r t        ||      || j                        S t        |      r5t        j                  dd|      } t        ||      || j                        S  t        ||      |      S )N_builderz__(.*)__z__r\1__)rM   r   r   resub)ry   method_namer_  r`  reverse_method_names        r2   _apply_binary_methodz"CodeGenerator._apply_binary_method  sn    S!,73,S4<<HHS!"$&&j+"N473 34S4<<PP(wsK(--r3   c                 ^   | j                  |j                        }| j                  |j                        }| j                  j	                  t        |j                              }|5| j                  |dj                  |j                  j                              | j                  |||      S )Nz8AST binary operator '{}' is not (currently) implemented.)r   leftright_method_name_for_bin_opr   rQ   r^  r   formatr`   rq  ry   rb   r_  r`  ro  s        r2   visit_BinOpzCodeGenerator.visit_BinOp  s    jj#jj$2266tDGG}E##D$^$e$efjfmfmfvfv$wy y((c3??r3   __add____sub____mul____truediv____floordiv____mod____pow__
__lshift__
__rshift____and____or____xor__ru  c                    | j                   j                  |       | j                  |j                         | j                   j	                         }| j
                  j                         }i }|j                  r| j                   j                  |       |j                         | _        i | _        | j                  |j                         | j
                  j                         }| j                   j	                         }g }g }g }	|D ]  }
|df|dffD ]V  \  }}|
|v s||
   j                  ||
   j                  k(  r+J d|
 d||
   j                   d| d||
   j                           |
|v s|
|v r|j                  |
       |j                  |
|v r||
   j                  n||
   j                         |	j                  |
|v r||
   j                  j                         n||
   j                  j                                |
|v r|
|vr||
   ||
<   |
|v s|
|vs||
   ||
<   " |j                         |j                         z  D ]  }
|
|v r||
   j                  }||
   j                  }||k(  sJ d|
 d| d	| d
       |j                  |
       |j                  |       |	j                  ||
   j                  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   r=  r  r   r   r   r}   r   r|   rQ   r9  r$  get_typekeys)ry   rb   r~   
then_block
else_block	then_defs	else_defsrX  r*  ir_ret_typesr@   defs
block_namethen_tyelse_tys                  r2   visit_then_else_blocksz$CodeGenerator.visit_then_else_blocks7  s   11*=%%dii0\\557
OO((*		;;LL55jA!,,.DK DO))$++6,,.I99;J 	 	0D&/%89f:M$N X j4<:??gdm.@.@@ X-dV=ASAS@T U##-,.Ed4jooEVXX@X
 y DI$5T"  9J4!5!5PYZ^P_PdPde##$-J.IdO$:$:$C$C$E3<T?3I3I3R3R3TV y T%:")$-	$y T%:")$-	$#	0( NN$y~~'77 
	CDu}o**Go**Gg% .&tf,A' K##*)1..% LLW%	$ 6 6 ? ? AB
	C )ZUI|[[r3   c           	         d}t        |       5 }|\  }}| j                  j                         }| j                  j                         }| j                  j                         }	| j                  j                  |       | j                  j	                  |j
                  ||       | j                  ||||      \  }
}}}}}}| j                  j                  |       |j                         r"|j                         rd}|	j                          |j                         s9|r7| j                  j                  |	|D cg c]  }|
|   j
                   c}       | j                  j                  |       |j                         s9|r7| j                  j                  |	|D cg c]  }||   j
                   c}       |r|D ]  }|	j                  |        d d d        |rs| j                  j                  	       t              D ]I  \  }}t        j                  j!                  |	j#                  |      |         }| j%                  ||       K y y c c}w c c}w # 1 sw Y   xY wr   )rv   r   create_blockr?  create_cond_branchr$  r  
has_returnerasehas_terminatorcreate_branchadd_argumentr=  r^   r   r"  r   rf   r  )ry   condrb   has_endif_blocksrr~   ip_blockr  r  endif_blockr  r  rX  r*  r  nr/   r   r@   
new_tensors                       r2   visit_if_top_levelz CodeGenerator.visit_if_top_leveln  s   d# 	1r "GX224J224J,,335KLL33H=LL++DKKZP ++D':zR YIy*j%L LL33J?$$&:+@+@+B"'!!#,,.?**;V[8\QR19L9L8\]LL33J?,,.?**;V[8\QR19L9L8\]& 1B,,R011	14 LL55kB$U+ 14%]]11+//!2DiPQlS
tZ01	  9] 9]-	1 	1s+   D+I9IAII3!I
II%c           	      &   t        |       5 }|\  }}| j                         \  }}| j                  j                         }|j                  r| j                  j                         nd }	| j                  ||||	      \  }
}}}	}}}| j                  ||       | j                  j                  |D cg c]  }|j                  | j                         c}|j                  d      }|j                  |j                                | j                  j                  |j                                t        |      dkD  r6| j                  j                  |D cg c]  }|
|   j                   c}       |j                  s|j                         }	n|	j                  |j                                | j                  j                  |j                                t        |      dkD  r6| j                  j                  |D cg c]  }||   j                   c}       d d d        t!              D ]I  \  }}t"        j$                  j'                  j)                  |      |         }| j+                  ||       K y c c}w c c}w c c}w # 1 sw Y   pxY w)NTr   )rv   r  r   r  r   r  r  create_if_opr6  r$  merge_block_beforeget_then_blockr?  rV  create_yield_opget_else_blockr^   r   r"  r   
get_resultr  )ry   r  rb   r  r~   r   r  last_locr  r  r  r  rX  r*  r/   if_opr  r   r@   r  s                       r2   visit_if_scfzCodeGenerator.visit_if_scf  s   d# 	SrJGQ<<>LB224J8<224J++D':zR NIy*j%A --b(;LL--PY.Z"rxx/E.Z\`\g\gimnE))%*>*>*@ALL33E4H4H4JK5zA~,,5-Qail.A.A-QR;;"113
--e.B.B.DELL33E4H4H4JK5zA~,,5-Qail.A.A-QR)	S, !' 	-GAt!--e.>.>q.A9Q<PJNN4,	- /[ .R .R)	S 	Ss8   B"J."I8
BJI=
(BJ;J
J8JJc           	         | j                  |j                        }t        |      r|j                  t        j
                  | j                        }t        | j                        j                  |      }| j                  r|r| j                  |d      | j                  s|s| j                  ||       y | j                  ||       y t        |      }t        |      t        vrO| j                  |dj!                  dj#                  d t        D              t        |      j$                              |r| j'                  |j(                         y | j'                  |j*                         y )Nrk  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   4   K   | ]  }|j                     y wrJ   r`   .0r   s     r2   	<genexpr>z)CodeGenerator.visit_If.<locals>.<genexpr>       !G!**!G   )r   testrM   tor   int1r   r   r   r   r   r  r  rZ   rQ   _condition_typesrv  r+   r`   r  r   r   )ry   rb   r  contains_returns       r2   r   zCodeGenerator.visit_If  s%   zz$))$T"778==4<<7@D3DKK@FFtLO~~/'' PQ Q !!$-''d3'-DDz!11''krr		!G6F!GGT
++-. . --dii8--dkk:r3   c           	      t   | j                  |j                        }t        |      rc|j                  t        j
                  | j                        }t        |       5  | j                         \  }}| j                  j                         }| j                  j                  |       t        j                  j                  | j                  |j                        | j                        }| j                  j                         }| j                  j                         }| j                  j                  |       t        j                  j                  | j                  |j                        | j                        }| j                  j                         }| j!                  ||       |j"                  |j"                  k(  s!J d|j"                   d|j"                          |j"                  }	|	t        j$                  k7  r|	j'                  | j                        gng }
| j                  j)                  |
|j*                  d      }|j-                  |j/                                |
rO| j                  j1                  |j/                                | j                  j3                  |j*                  g       | j                  j1                  |j/                                |j-                  |j5                                |
rO| j                  j1                  |j5                                | j                  j3                  |j*                  g       |
r/t        j                  j7                  |j9                  d      |	      nd cd d d        S t;        |      }t#        |      t<        vrO| j?                  |djA                  djC                  d t<        D              t#        |      jD                              |r| j                  |j                        S | j                  |j                        S # 1 sw Y   y xY w)	Nrk  zAternary expression with dynamic condition has inconsistent types r   Tr   r  r  c              3   4   K   | ]  }|j                     y wrJ   r  r  s     r2   r  z,CodeGenerator.visit_IfExp.<locals>.<genexpr>  r  r  )#r   r  rM   r  r   r  r   rv   r  r  r=  r"  r#  r   r   r   r  rQ   r!  r6  r  r$  r  r  r?  r  r  r   r  rZ   r  r   rv  r+   r`   )ry   rb   r  r  r  r  then_valr  else_valr   ret_type_irr  s               r2   r   zCodeGenerator.visit_IfExp  s+   zz$))$T"778==4<<7@D!$' !d#@@BH!\\668
99*E#==33DJJtyy4I4<<X!\\==?
!\\668
99*E $==33DJJt{{4KT\\Z!\\==?
11"h?}}5 |WX`XeXeWffkltlylykz{|5#==@HHMM@Yx~~dll;<_a11+t{{DQ--e.B.B.DELL;;E<P<P<RSLL00(//1BC778L8L8NO--e.B.B.DELL;;E<P<P<RSLL00(//1BCNYx}}++E,<,<Q,?J_cC!d !dF (-D Dz!11''krr		!G6F!GGT
++-. . zz$)),,zz$++..]!d !ds   L#P..P7c                      y rJ   r   r   s     r2   
visit_PasszCodeGenerator.visit_Pass  s    r3   c                    t        |j                        dk(  rt        |j                        dk(  s| j                  |d      | j	                  |j
                        }| 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                  j                  t        |j                  d               }|8| j                  |dj                  |j                  d   j                              | j!                  |||      S )Nr   z1simultaneous multiple comparison is not supportedr   z<AST comparison operator '{}' is not (currently) implemented.)rV  comparatorsopsr   r   rs  rZ   rQ   r   Isr   IsNot_method_name_for_comp_opr   rv  r`   rq  )ry   rb   r_  r`  	lhs_value	rhs_valuero  s          r2   visit_ComparezCodeGenerator.visit_Compare  s@   D$$%*s488}/A##D*]^^jj#jj))!,-(-	(-	&Y)344		)Yi7883377TXXa[8IJ##T[[\`\d\def\g\p\pqs s((c3??r3   __eq____ne____lt____le____gt____ge__r  c           
         | j                  |j                        }| j                  j                  t	        |j
                              }|*| j                  |d|j
                  j                   d      t        |      r t        ||      | j                        S 	  t        ||             S # t        $ r, | j                  |d| dt	        |      j                         w xY w)NzAST unary operator 'z!' is not (currently) implemented.rk  z)' is not (currently) implemented on type )r   operand_method_name_for_unary_opr   rQ   r^  r   r`   rM   r   r   AttributeError)ry   rb   r  rc   s       r2   visit_UnaryOpzCodeGenerator.visit_UnaryOp  s    **T\\*++//TWW>:##D,@AQAQ@RRs*tuuW%'77B'>>	t'77B')) 	t##,RD0YZ^_fZgZpZpYqrt t	ts   B0 05C%__neg____pos____not__
__invert__r  c           
      B   t        |       5 }|\  }}| j                         \  }}| j                  j                         }| j                  j	                  |       | j
                  j                  |       | j                  |j                         | j
                  j                          | j                  }|j                          g }	g }
g }|D ]  }||v st        ||         sJ d| d       t        ||         sJ d| d       ||   j                  ||   j                  k(  s+J d| d||   j                   d||   j                   d       |	j                  |       |
j                  ||   j                         |j                  ||           | j                  ||       | j                  j                  |
D cg c]  }|j!                  | j                         c}|D cg c]  }|j"                   c}      }| j                  j%                  |j'                         |
D cg c]  }|j!                  | j                         c}      }| j                  j	                  |       t)        |	      D ]`  \  }}t*        j,                  j/                  |j1                  |      |
|         | j2                  |<   | j2                  |   | j                  |<   b | j5                  |j6                        }| j                  j9                  |       | j                  j;                  |j"                  t=        t?        |            D cg c]  }|j1                  |       c}       | j                  j%                  |jA                         |
D cg c]  }|j!                  | j                         c}      }| j                  j	                  |       t)        |	      D ]`  \  }}t*        j,                  j/                  |j1                  |      |
|         | j2                  |<   | j2                  |   | j                  |<   b | j
                  j                  |       | j                  |j                         | j
                  j                          | j                  }g }|D ]  }||v s|j                  ||           | j                  jC                  |D cg c]  }|j"                   c}       d d d        t)        	      D ]U  \  }}t*        j,                  j/                  jE                  |      
|         }|| j2                  |<   || j                  |<   W |jF                  D ]  }J d        y c c}w c c}w c c}w c c}w c c}w c c}w # 1 sw Y   xY w)	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.zNot implemented)'rv   r  r   r  r=  r   r9  r  r   popr   r  rM   rQ   r  create_while_opr6  r$  create_block_with_parent
get_beforer^   r   r"  r   rf   r|   r   r  r?  create_condition_oprangerV  	get_afterr  r  r   r   r  r   )ry   rb   r  r~   r   r  r  dummy	loop_defsrX  r*  	init_argsr@   r/   rf   while_opbefore_blockr   r  after_blockyieldsynew_defr  s                           r2   visit_WhilezCodeGenerator.visit_While'  s    d# @	Er$&!G\<<>LB LL--/ELL55e<NN!!$'))$))4NN IKKM EII! 47?,Yt_=mAZ[_Z``l?mm=,WT];k?XY]X^^j=kk;$T?//74=3E3EE L06HQUI[I[H\ ]11:41E1E0F GKLLE LL&$$Yt_%9%9:$$WT]34 --b(;||33V_4`PRRXXdll5K4`KT4UCSZZ4UWH  <<@@ATATAVclAm]_"((4<<BXAmoLLL55lC$U+ :4$,MM$8$89I9I!9LiXYl$[D!(,D(9%: ::dii(DLL33LALL,,T[[X]^abk^lXm:nST<;K;KA;N:no,,??@R@R@Tbk@l\^$,,AW@lnK LL55kB$U+ :4$,MM$8$89KYWX\$ZD!(,D(9%: NN!!$'))$))4NN IF! 37?MM)D/23 LL((F)Cq!(()CDA@	EF !' 	,GAtmm**8+>+>q+A9Q<PG 'DKK$+DOOD!	,
 KK 	6D+++5	6M 5a4U Bn ;o Am *DA@	E @	Esi   C	VC,V"U7
$V*U<=/V,"VC<V
V"/V"V3C7V+/VV
-V7VVc                    |j                   j                  j                  dk(  sJ | j                  |j                        }| j                  |j
                        }t        |      r|j                  || j                        S ||   S )Nr\  rk  )	r   	__class__r`   r   rY   slicerM   __getitem__r   )ry   rb   r_  slicess       r2   visit_SubscriptzCodeGenerator.visit_Subscriptt  sm    xx!!**f444jj$DJJ'S!??6DLL?AA6{r3   c                 ^    |j                   D cg c]  }| j                  |       c}S c c}w rJ   )dimsr   )ry   rb   dims      r2   visit_ExtSlicezCodeGenerator.visit_ExtSlice|  s"    +/995C

3555s   *c                       j                  |j                  j                        }|j                  j                  D cg c]  } j                  |       }}t	         fd|j                  j
                  D              }|t        j                  k(  r ||i |}t        |j                  j                  |j                  j                  |j                  j                        }|D ]z  }t        |       j                  |j                  j                   <    j#                  |j$                         |j&                  D ]"  }	t(        j*                  j-                   |	       $ | y d }
|t        j                  u r9 ||i |}|j                  }|j                  }|j                  }|j.                  }
n|t        u rt1        |      dkD  r|d   n# j                  t)        j2                  d            }t1        |      dkD  r|d   n' j                  |j                  j                  d         }t1        |      dkD  r|d   n# j                  t)        j2                  d            }nt5        d      d}t7        |      r+|j                  dk  rt        |j                         }d}||}}t        j8                  j;                  | j<                        }t        j8                  j;                  | j<                        }t        j8                  j;                  | j<                        }|j>                  jA                         r4|j>                  jA                         r|j>                  jA                         s3tC        d|j>                   d	|j>                   d	|j>                   d
      t        jD                  jG                  |j>                  |j>                        }t        jD                  jG                  ||j>                        }|jI                   j<                        }|jJ                  t        j8                  j>                  jL                  jN                  k(  }|jP                  }|jP                  }|jP                  } j<                  jS                  |||      } j<                  jS                  |||      } j<                  jS                  |||      } j<                  jU                  |      } jW                  |j                  j                   t        j8                  jY                  ||             t[               5 }|\  }} j]                         \  }} j<                  j_                         } j<                  ja                  |        jb                  je                  |        j#                  |j$                          jb                  jg                          |ji                          g }g }g } jj                  D ]+  }||v s	tm         jj                  |         s
J | d       tm        ||         sJ  jj                  |   jn                  ||   jn                  k(  s5J d| d||   jn                   d jj                  |   jn                   d       |je                  |       |je                  t        j8                  j;                  ||    j<                               |je                  t        j8                  j;                   jj                  |    j<                               .  jq                  ||        j<                  js                  ||||D cg c]  }|jP                   c}      }|
+|ju                  d j<                  jw                  |
              jb                  je                  |        j<                  ja                  |jy                  d             |j{                          _        i  _5        t}        |      D ]c  \  }} jW                  |t        j8                  jY                  |jy                  d      j                  |dz         ||   jn                               e  j#                  |j$                          jb                  jg                          g } jj                  D ]M  }||v s|je                  t        j8                  j;                   jj                  |    j<                               O t1        |      dkD  r3 j<                  j                  |D cg c]  }|jP                   c}       |jy                  d      j                         }|j                         dk(  sJ d        j<                  ja                  |jy                  d             |j                         }|r8 j<                  j                  ||      } j<                  j                  ||      } j                  |j                  j                      jP                  j                  |        jW                  |j                  j                   t        j8                  jY                  ||             d d d        t}              D ]Q  \  }} jW                  |t        j8                  jY                  j                  |      |   jn                               S |j&                  D ]  }	J d        y c c}w c c}w c c}w # 1 sw Y   xY w)Nc              3   @   K   | ]  }j                  |        y wrJ   r   r  keywordry   s     r2   r  z*CodeGenerator.visit_For.<locals>.<genexpr>  s     Q74::g.Q   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  ztt.num_stagesz7We use SCF, so the loop body should only have one blockz)Don't know what to do with else after for)Hr   iterr   rd   r   keywordsr   static_ranger  startrY   endstepr   r|   r/  r   r  r   r   r   r  r   
num_stagesrV  NumRuntimeErrorrO   r"  r#  r   r"   r!   r%  semanticinteger_promote_implr6  r%   r#   r$   r$  create_int_castcreate_undefr  r   rv   r  r  r=  r   r9  r  r  r   rM   rQ   r  create_for_opset_attrget_int32_attrget_bodyr}   r^   rf   r  
get_parentsizeget_induction_var
create_sub
create_addreplace_all_uses_withr  ) ry   rb   IteratorClassrf   	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  rX  r@   for_opr  for_op_regions    `                               r2   	visit_ForzCodeGenerator.visit_For  s   

499>>204		?TZZ_?	?Qdii>P>PQQH111$i?;?H !5!5x||7I7I8==K^K^_L! >.7lDKKNN+--dii8 KK >DOO11$=>>
 
HNN*$i?;?H BB==D!,,Je# "%Y!!31CGGAJ9OB!$Y!!31DIINNSTDU9VB#&y>A#59Q<4::cggaj;QDbcc4::>djj[)D MB]]%%b$,,7]]%%b$,,7}}''dll;xx (9ARARATNrxxjXZ[][c[cZddfgkgqgqfrrstuu##88288L##88$**M]]4<<0
--1D1D1O1O1V1VVYYYY{{\\))"j,G\\))"j,G||++D*lK\\&&z2t{{~~x}}';';B'HId# @	Nr$&!G\<<>LB LL--/ELL55e<NN!!$'))$))4NN KKM IFE a7?,T__T-BC\v^E\\C,WT];;;??40559K9KK L06HQUI[I[H\ ]1151F1K1K0L MKLLK
 LL&$$X]]%=%=gdmT\\%Z[MM(--":":4??4;PRVR^R^"_`a --b(;\\//BU^>_cszz>_`F%1L1LZ1XYNN!!$'LL55fooa6HI!,,.DK DO$U+ j4tX]]%9%9&//!:L:P:PQRUVQV:WY_`aYbYgYg%hij))$))4NN F a7?MM(--":":4??4;PRVR^R^"_`a
 6{Q,,-G1ahh-GH"OOA.99;M %%'1,g.gg, LL55fooa6HI))+B\\,,R4\\,,R4KK'..DDRHNN4;;>>8==+?+?G+LMA@	NF !' 	]GAtNN4!5!5f6G6G6JFSTINN![\	] KK 	6DEEE5	6A @v ?`( .Hk@	N @	Ns@   k5?ClElk:/El A/l/k?
D4l:
llc                     | j                  |j                        }| j                  |j                        }| j                  |j                        }t	        |||      S rJ   )r   lowerupperr  r  )ry   rb   r)  r*  r  s        r2   visit_SlicezCodeGenerator.visit_Slice  sG    

4::&

4::&zz$))$UE4((r3   c                 8    | j                  |j                        S rJ   )r   rY   r   s     r2   visit_IndexzCodeGenerator.visit_Index  s    zz$**%%r3   c                 P    |j                   | j                  |j                        fS rJ   )rf   r   rY   r   s     r2   visit_keywordzCodeGenerator.visit_keyword  s    xxDJJ///r3   c                     | j                   sy | j                  |j                        }|j                  | j                  |j                        nd}t        j
                  j                  ||| j                        S )Nr   rk  )r   r   r  msgr   r"  device_assertr   )ry   rb   r  r1  s       r2   visit_AssertzCodeGenerator.visit_Assert  s[    zzzz$))$&*hh&:djj"}}**4t||*LLr3   rc   c                    t        j                  |j                  g|i |}|j                  D cg c]  }||   	 }}|D cg c]  }t	        |      r|n
t        |       }}t               }t        |      D cg c]  \  }}t        |      s| }}}|D ci c]  }|||   
 }	}t        |      D cg c]  \  }}||v rd n| }}}|D cg c]  }||j                   }
}|D cg c]  }||j                   }}t        |j                  ||	      }| j                  j                  |      st        j                   g |      }|j"                  }t%        |      \  }}|j&                  | j&                  n|j&                  }t)        | j*                  ||||	| j                  ||| j,                  |j.                  ||| j0                  j2                  | j0                  j4                  |      }	 |j7                  |j9                                |jD                  }|| j,                  |<   n| j,                  |   }| j                  jG                  |      }| j0                  jI                  ||
      }|jK                         dk(  s|y |jK                         dk(  rtM        |jO                  d      |      S g }tQ        |jK                               D ]/  }|jS                  tM        |jO                  |      ||                1 tU        |      S c c}w c c}w c c}}w c c}w c c}}w c c}w c c}w # t:        $ r1}t=        | j>                  j@                  | jB                  d       |d }~ww xY w)N)
r   r   r   r   r]   rp   rr   r   r   r   r   r   )+rk   getcallargsrc   ra   rM   r   r   r^   rO   r$  rQ   rF   r`   r   has_functionr   function_type__globals__rt   r   r   r   r   r]   r   r   r   r   r   	Exceptionr   r   r_   r   r   get_functioncallget_num_resultsr   r  r  r9  rV   )ry   rc   rd   r   r@   rf   r   r   
constexprsrB   arg_vals	arg_typesfn_namer   r   rp   rr   r   rx   ecallee_ret_typesymbolcall_opresultss                           r2   call_JitFunctionzCodeGenerator.call_JitFunction  s   ""255:4:6:')||4tT
44MQRc(-9S>ARRV
&/oLFAss9KaL
L)34AQQZ4	4?HOVQZS0OO*.B3#/CJJBB)-A#SXXA	ABKKI>{{''0 ..r9=I^^F$5b$9!Iz"$(("2DJJE%dllIvzS\eiepep-/wW[WnWn/1{{i\f.2ll.B.BPTP\P\PhPhpuwIT
+
 (00O/>D##G,"55g>O))'2,,##FH5""$)_-D$$&!+',,Q/AA G72245 Rvg&8&8&;_Q=OPQR>!W 5R M4OBA  T&t{{tLRSSTsS   LL=LLL$9L)L/L/.L46L47L9 9	M3,M..M3c                     t         j                  |j                              } j                  j	                  |      }|	 | |      S t         fd|j                  D              }|j                  D cg c]  } j                  |       }}|t        j                  j                  u r j                  sy t        |t              r t        |||        j                  |||      S t!        |d      rt#        |j$                        st        j                  j'                  |      rIt         j(                        }t+        j,                  |      }d|j.                  v r |d<   	  ||i ||S | j8                  j;                         v rt=        t         |      } ||i |S c c}w # t0        $ r'}	t3         j4                  j6                  |d       |	d }	~	ww xY w)Nc              3   @   K   | ]  }j                  |        y wrJ   r   r  s     r2   r  z+CodeGenerator.visit_Call.<locals>.<genexpr>N  s     D74::g&Dr  __self__rk  
_generator)rZ   r   r    statically_implemented_functionsr   r   r  rd   r   r"  r2  r   rK   r   rg   rF  hasattrrM   rI  
is_builtinr   rk   	signature
parametersr9  r   r   r_   r   rY  r,   )
ry   rb   rc   static_implementationkwsrf   rd   extra_kwargssigrA  s
   `         r2   r   zCodeGenerator.visit_CallH  s   !$**TYY"78 $ E E I I" M ,(t44DdmmDD+/995C

355,,,::b+&4T*((T377B
#(9"++(F8==KcKcdfKg6L##B'Cs~~--1\*	K47<7377 ''..00+T2D433 6  K 't{{dCJKs   8F(*	F- -	G6"GGc                 ,    t        |j                        S rJ   r   rY   r   s     r2   visit_ConstantzCodeGenerator.visit_Constantj  s    $$r3   rb   c                    t        |j                        dk7  r| j                  |d      | j                  |j                  d         }| j                  |j                  d         }| j                  j                  t        |j                              }|5| j                  |dj                  |j                  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.)rV  rY  r   r   _method_name_for_bool_opr   rQ   r^  rv  r`   rq  rw  s        r2   visit_BoolOpzCodeGenerator.visit_BoolOpm  s    t{{q ##vx xjjQ(jjQ(3377TWWF##QXXY]Y`Y`YiYijl l((c3??r3   logical_and
logical_orrX  )      c                 ,    t        |j                        S rJ   rU  r   s     r2   visit_NameConstantz CodeGenerator.visit_NameConstant}  s    TZZ((r3   c                 ,    t        |j                        S rJ   )r   r  r   s     r2   	visit_NumzCodeGenerator.visit_Num  s    TVV$$r3   c                 >    t        t        j                  |            S rJ   )r   r   literal_evalr   s     r2   	visit_StrzCodeGenerator.visit_Str  s    S--d344r3   c                     | j                  |j                        }t        |      r;|j                  dk(  r,t        j
                  j                  |d| j                        S t        ||j                        S )NT)r   r   )r   )	r   rY   rM   r   r   r  permuter   r   )ry   rb   r_  s      r2   r   zCodeGenerator.visit_Attribute  s[    jj$S!yyC((00fdll0SSsDII&&r3   c                 D    t         j                  j                  | |       y rJ   r  r   s     r2   
visit_ExprzCodeGenerator.visit_Expr  r  r3   c                      y rJ   r   r   s     r2   visit_NoneTypezCodeGenerator.visit_NoneType  r   r3   c           
      p   t        |j                        }t        |      D ]  \  }}t        |t        j
                        rt        |j                        ||<   :t        |t        j                        r|j                  }| j                  |j                        }t        |      s'| j                  |dt        t        |            z         |dk  rdndt        |      z   dz   j                  |j                        ||<   t!        dj                  t        |                   dj#                  |      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   )rU   rY  r^   rK   r   Constantr&   rY   FormattedValue
conversionr   rO   r   rQ   chrrv  AssertionErrorr+   )ry   rb   rY  r   rY   conversion_code	evaluateds          r2   visit_JoinedStrzCodeGenerator.visit_JoinedStr  s   dkk"!&) 	wHAu%.,q	E3#5#56"'"2"2 JJu{{3	$Y/++xd9o./0 0 &5q%8TdSEY>Y\_>_gghqhwhwxq	$%a%h%himnsit%uvv	w wwvr3   c           	         |y t        j                         5  t        j                  dt               t        j                  dt               | j
                  }| j                  j                         }|| _        t        |d      rnt        |d      rb| j                  j                  | j                  | j                  |j                  z   |j                         | j                  j                         }	 t        | =  |      }|r"|| _        | j                  j                  |       |cd d d        S # t         $ r  t"        $ r:}t!        | j$                  j&                  | j
                  t)        |            d d }~ww xY w# 1 sw Y   y xY w)Nignorelineno
col_offset)warningscatch_warningssimplefilterDeprecationWarningPendingDeprecationWarningr   r   r  rL  r   rp   rr   rx  ry  superr   r   r9  r   r_   r>   )ry   rb   	last_noder  rE   rA  r  s         r2   r   zCodeGenerator.visit  sB   <$$& 	 !!(,>?!!(,EFI||++-H DMtX&74+F$$T^^T__t{{5RTXTcTcd<<//1ZgmD)  )$$X.1	 	 $  Z 't{{tAwOUYYZ	 	s0   CF6D4%F4F 5E;;F  FFc                 j    | j                  |dj                  t        |      j                              )Nzunsupported AST node type: {})r   rv  rQ   r`   r   s     r2   r   zCodeGenerator.generic_visit  s,    &E&L&LTRVZM`M`&abbr3   c                    t        |j                        }d|cxk  rdk  r"n t        d      t        |j                        rt        d      t	        | j                  |j                  d               }t        |t              st        d      |sQ|dk(  rd}n	 | j                  |j                  d         }t        | j                  j                  |t	        |            y # t        $ r}dt        |      z   dz   }Y d }~Jd }~ww xY w)	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: >)rV  rd   r  r%  rZ   r   rK   r   NotImplementedErrorr9  r>   r   r   r_   )ry   rb   	arg_countpassedr   rA  s         r2   execute_static_assertz#CodeGenerator.execute_static_assert  s    		N	I""[\\ (+4=='9[\\%djj1&>?&$'% D  A~X"jj16G .dkkootEYZaEbcc	 ! XG$q'QTWWGXs   C$ $	D-DDc                 4     dt         j                  f fd}|S )Nrb   c                      fd|j                   D        D ci c]  \  }}|t        |       }}}|j                  D cg c]  }t         j                  |             }}t	         |i |      S c c}}w c c}w )Nc              3   @   K   | ]  }j                  |        y wrJ   r   r  s     r2   r  z=CodeGenerator.static_executor.<locals>.ret.<locals>.<genexpr>  s     #UGDJJw$7#Ur  )r  rZ   rd   r   r   )ry   rb   r@   rY   rQ  rf   rd   	python_fns   `      r2   rE   z*CodeGenerator.static_executor.<locals>.ret  s     $Vt}}#UD% *511C  FJYYOc(C9ODOY4455 Ps   A6!A<)r   r   )r  rE   s   ` r2   static_executorzCodeGenerator.static_executor  s    	6CHH 	6 
r3   rK  )NNFNFNr   )~r`   r   r   r   r   r   r&   rz   rV  rU   r  floatintrK   r   r   r   r   updater   r"  device_printminimummaximumr   r   r   r	   r   r   r  r  r  r  r   r  r   r   rO  rQ  rT  r   r   r   rd  rf  ri  rq  rx  r   AddSubMultDivFloorDivModPowLShiftRShiftBitAndBitOrBitXorru  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/  r3  rF  r   rV  BoolOprY  AndOrrX  boolopsysversion_infor_  ra  rd  r   ri  rk  ru  r   r   r   r  r  static_assertstatic_printr   rK  r   r   __classcell__)r  r   r  s   00@r2   r   r      s    jnMN$0al $0X`aeXf$0,4SM$0L BEdESXZ]_ikr@s(t1Q(ttCH~t	(--,,-	  !	  ! L,\&c &%	0A*B &t &"
2X4AF&'"(,+.
22.@ 	)n

L

L

I		8

I>T$s||"4c"9: 5\n"1J-6;63/j@$ 	#))Xsvvx(TWTZTZ\dfifmfmow<d4		?C#78 t 	)SXXy#''9cjjR^?tD$5s$:; K6Z6D6L)&0U38_ 0MC M-"; -"^  D%
@ 
@ >AWWmUXU[U[]i<jd4

#3S#89j
& 	)	%	5'2$:c#(( t ,
 	##%:""OE$:_S!_S!	Q$d68SXXJO3L+L&M W )us   Mr   c                     d}t        |       D ]:  \  }}|t        |      z  }||j                  v r|dz  }||j                  v s6|dz  }< |S )Nr   r5   d)r^   r&   
equal_to_1divisible_by_16)rN  specializationsuffixr   r   s        r2   kernel_suffixr    sd     F)$ 1#a&)))cMF...cMF Mr3   c                     |j                   } fd}|j                  j                         D ci c]  \  }} ||      | }	}} j                  j	                         }
 j                  |      }t        |j                  j                               }|j                  D ci c]  }|||v r
||   dk(  rdnd }}|j                  D ci c]  }|dg }}|	j	                         }|j                  |       |j                  j                         D cg c]  \  }}||j                  vst        |      ! }}}t               \  }}t        j                  g |      }t!        |||
|| |d||||      }|j#                   j%                                |j&                  }||_        |S c c}}w c c}w c c}w c c}}w )Nc                 ^    t        | t              rj                  j                  |       S | S rJ   )rK   r&   ra   index)r   rc   s    r2   <lambda>zast_to_ttir.<locals>.<lambda>  s#    As1C**1-  r3   i1Tr   )ztt.divisibility   )
r   rB   r   r   r   r   rp   rr   r   r   )attrsrB   itemsr8  r}   r>   rU   rN  rY  r  r  r  r   rt   r   r7  r   r   r   r   r   )rc   r  r   r   r   r  cst_keykeyrY   rB   r   r   tysknew_constants	new_attrsall_constantsr(  r?  rp   rr   r   rx   rE   s   `                       r2   ast_to_ttirr    s     EJG7E7O7O7U7U7WXeu$XIX^^  "FGGN+M
~''..0
1CLQL\L\]qQSSVt^B]M]7<7L7LM!,--MIMNN$M'*8*B*B*H*H*Jp$!QaWeWoWoNo1pIp-b1Iz&&r95Igy=hu%'IYb)3WR]_I OOBHHJ


CCKJ- Y
 ^M qs   F+"F1
F6F;,F;)3r   rk   rm  r  rz  r  r  typingr   r   r   r   r   r   r	   r   r   _C.libtritonr   r   r   r   runtime.jitr   runtimer   errorsr   r   r   typesr   r   rF   r   rM   rO   rS   rW   rZ   rg   rt   r  rQ   r  rv   r  r   r   r  r  r   r3   r2   <module>r     s    
  	 
  	  D D D   3 3 ' ! a a %$	! ! !$S $T $Q Q Q(S (T (6C 6!$ #tDz* 3 3*X%COO X%vXCOO Xv 
r3   