
    ci                       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Zd dlZd dl	m
Z
 d dlmZ d dlmZmZmZmZmZmZmZmZmZ ddlmZmZ ddlmZmZ ddlmZmZmZm Z! dd	l"m#Z#m$Z$m%Z% dd
l&m'Z'm(Z(m)Z)m*Z*m+Z+m,Z, ddl-m.Z.m/Z/m0Z0 ddl1m2Z2m3Z3m4Z4 d Z5d Z6dede7fdZ8dede7fdZ9dede7fdZ:dede7fdZ;dede7fdZ<d Z=d Z>d Z?dee$   fdZ@deej                     dee%   fdZBe7eC eDd      hZEd ZFd  ZG G d! d"      ZH G d# d$e j                        ZJ G d% d&      ZK e
d'(       G d) d*             ZL G d+ d,e j                        ZMd.d-ZNy)/    N)	dataclass)
ModuleType)	AnyCallableDictOptionalTupleTypeUnionIterableList   )knobslanguage)irgluon_ir)	constexpr	str_to_tytensortuple)_unwrap_if_constexpr
base_value	base_type)get_jit_fn_file_lineget_full_nameJITCallableBoundConstexprFunctionConstexprFunctionJITFunction)find_paths_ifget_iterable_pathset_iterable_path   )CompilationErrorCompileTimeAssertionFailureUnsupportedLanguageConstructc                 Z    d}t        j                  ||       st        d| d|  |       | S )Nz^[a-zA-Z_][a-zA-Z0-9_]*$zinvalid z identifier: )rematchr$   )nametypepatterns      X/var/www/html/engine/venv/lib/python3.12/site-packages/triton/compiler/code_generator.pycheck_identifier_legalityr.      s4    )G88GT"$}TFCTJJK    c                    dj                  |D cg c]  }|j                          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	| }|||j                         z  }|S c c}w c c}w )
N_c._d_'_sq_[]__)joinmanglesortedreprreplace)	r*   arg_tys	constantscaller_contexttymangled_arg_namesimangled_constantsrets	            r-   	mangle_fnrG       s    !@""))+!@AviGX!Y!QCqil);(<"=!YZ)11#u=)11#v>)11#s;CCCMF"&'r*;)<
=C!~$$&&J "A!Ys   CCoreturnc                 "    t        | t              S N)
isinstancer   rH   s    r-   _is_triton_valuerN   .   s    a$$r/   c                 "    t        | t              S rK   )rL   r   rM   s    r-   _is_triton_tensorrP   2   s    a  r/   c                 l    | d u xs/ t        | t        t        j                  j                  t
        f      S rK   )rL   r   r   coredtyper   rM   s    r-   _is_constexprrT   6   s)    9T
1y(--2E2E{&STTr/   c                     t        |       xr5 | j                  j                         xr | j                  j                  dk7  S Nr#   )rP   r+   is_blocknumelrM   s    r-   _is_non_scalar_tensorrY   :   s/    QMQVV__%6%L166<<1;LMr/   c                 .    t        | t        t        f      S rK   )rL   listr   rM   s    r-   _is_list_liker\   >   s    a$''r/   c                     |j                   rbt        |      D ]S  \  }}t        |      rt        |      st	        |j
                  | d|j                   d|j                  |    d|        y y )Nz	Function z= is marked noinline, but was called with non-scalar argument :)noinline	enumeraterT   rY   r&   src__name__	arg_names)nodefnargsidxargs        r-   _check_fn_argsri   B   s    	{{!$ 	HC %*?*D2FFD},ijljvjvwzj{i||}  B  ~C  D 	 r/   c                 b    t        | t              xr t        | t              xr t	        | d      S )N_fields)rL   r+   
issubclassr   hasattr)vals    r-   _is_namedtuplero   L   s'    c4 WZU%;WY@WWr/   c                    t        t        |             r| j                  }nDt        | t        j
                        r| j                  j                  }nJ dt        |               | D cg c]
  } ||       }}|D cg c]  }|t        |      n| }}|D cg c]  }|j                   }}t	        j
                  |t	        j                  ||            S c c}w c c}w c c}w )NzUnsupported type )	ro   r+   rk   rL   r   r   fieldsr   
tuple_type)valuere   rq   vvalstypess         r-   _apply_to_tuple_valuesrw   P   s    d5k"	E8>>	*""7)$u+77u !aBqE!D!6:;AIIaL1,;D;!"QVV"E">>$ 3 3E6 BCC ";"s   *C?CC$valuesc                 :    g }| D ]  }|j                  |        |S rK   )_flatten_ir)rx   handlesrt   s      r-   flatten_values_to_irr|   ^   s'    G 	gNr/   r{   rv   c              #   p   K   d}|D ]  }|j                  | |      \  }}|  |t        |       k(  sJ y wNr   )_unflatten_irlen)r{   rv   cursorrB   rs   s        r-   unflatten_ir_valuesr   e   sI     F ((&9v S\!!!s   46c                 j    g }| j                  |       | j                  j                  |d      \  }}|S r~   )rz   r+   r   )rn   r{   cloner1   s       r-   _clone_triton_valuer   p   s2    GOOGxx%%gq1HE1Lr/   c                     | j                         D ci c]  \  }}|t        |      rt        |      n|  c}}S c c}}w rK   )itemsrN   r   )scoper*   rn   s      r-   _clone_scoper   w   s;    ^c^i^i^klQZQUWZD.>s.C%c*Lllls   #;c                       e Zd Zd Zd Zd Zy)enter_sub_regionc                     || _         y rK   )	generator)selfr   s     r-   __init__zenter_sub_region.__init__}   s	    "r/   c                    t        | j                  j                        | _        t        | j                  j                        | _        i | j                  _        | j                  j                  j                         | _        | j                  j                  j                         | _
        | j                  | j                  fS rK   )r   r   lscopeliveins
local_defs	prev_defsbuilderget_insertion_blockinsert_blockget_insertion_pointinsert_point)r   s    r-   	__enter__zenter_sub_region.__enter__   s    #DNN$9$9:%dnn&?&?@$&! NN22FFH NN22FFH||T....r/   c                     | j                   j                  j                  | j                         | j                  | j                   _        | j                  | j                   _        y rK   )r   r   restore_insertion_pointr   r   r   r   r   )r   rf   kwargss      r-   __exit__zenter_sub_region.__exit__   s@    66t7H7HI $$(NN!r/   N)rb   
__module____qualname__r   r   r    r/   r-   r   r   {   s    #/3r/   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j2                  defdZdej6                  defdZy)ContainsReturnCheckerc                     || _         y rK   )gscope)r   r   s     r-   r   zContainsReturnChecker.__init__   s	    r/   rI   c                 ,     t         fd|D              S )Nc              3   @   K   | ]  }j                  |        y wrK   visit).0sr   s     r-   	<genexpr>z5ContainsReturnChecker._visit_stmts.<locals>.<genexpr>   s     /Q4::a=/   )any)r   bodys   ` r-   _visit_stmtsz"ContainsReturnChecker._visit_stmts   s    /$///r/   c                      yNFr   )r   re   s     r-   _visit_functionz%ContainsReturnChecker._visit_function   s     r/   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_fieldsrL   r[   ASTr   )r   rd   rF   r1   rs   items         r-   generic_visitz#ContainsReturnChecker.generic_visit   s    - 	/HAu%&! 6D!$0!5TZZ%56 E377+.TZZ.	/ 
r/   rd   c                 Z   t        |j                  t        j                        rm|j                  j                  | j
                  v rJ| j
                  |j                  j                     }t        ||j                        }| j                  |      S y| j                  |j                        S r   )
rL   rs   r   Nameidr   getattrattrr   r   )r   rd   rs   re   s       r-   visit_Attributez%ContainsReturnChecker.visit_Attribute   sw     djj#((+zz}}+DJJMM2UDII.++B//zz$**%%r/   c                     t        |j                        t        j                  u ry|j                  | j
                  v r*| j
                  |j                     }| j                  |      S yr   )r+   ctxr   Storer   r   r   )r   rd   re   s      r-   
visit_Namez ContainsReturnChecker.visit_Name   sO    >SYY&77dkk!TWW%B''++r/   c                      yNTr   r   rd   s     r-   visit_Returnz"ContainsReturnChecker.visit_Return       r/   c                      yr   r   r   s     r-   visit_Assignz"ContainsReturnChecker.visit_Assign        r/   c                      yr   r   r   s     r-   visit_AugAssignz%ContainsReturnChecker.visit_AugAssign   r   r/   c                 8    | j                  |j                        S rK   r   r   r   s     r-   visit_Modulez"ContainsReturnChecker.visit_Module         ++r/   c                 8    | j                  |j                        S rK   r   r   s     r-   visit_FunctionDefz'ContainsReturnChecker.visit_FunctionDef   r   r/   c                     | j                  |j                        }|j                  r|xs | j                  |j                        }|S rK   )r   r   orelse)r   rd   rF   s      r-   visit_IfzContainsReturnChecker.visit_If   s=     		*;;7**4;;7C
r/   c                 r    | j                  |j                        xs | j                  |j                        S rK   )r   r   r   r   s     r-   visit_IfExpz!ContainsReturnChecker.visit_IfExp   s'    zz$))$?

4;;(??r/   c                 8    | j                  |j                        S rK   )r   funcr   s     r-   
visit_Callz ContainsReturnChecker.visit_Call   s    zz$))$$r/   N)rb   r   r   r   boolr   r   r   r   	Attributer   r   r   Returnr   Assignr   	AugAssignr   Moduler   FunctionDefr   Ifr   IfExpr   Callr   r   r/   r-   r   r      s   0D 0T 	T 	&CMM &d &sxx D     
CMM d 
, , ,,coo ,$ ,SVV  @		 @d @%sxx %D %r/   r   c                       e Zd Zd Zdej
                  dee   deej                     fdZ	dej
                  deej                     fdZ
dej
                  fdZd Zy	)
ASTFunctionc                 <    || _         || _        || _        || _        y rK   )	ret_types	arg_typesr@   attrs)r   r   r   r@   r   s        r-   r   zASTFunction.__init__   s    """
r/   r   rv   rI   c                 B    g }|D ]  }||j                  ||        |S rK   )_flatten_ir_types)r   r   rv   ir_typesrB   s        r-   flatten_ir_typeszASTFunction.flatten_ir_types   s5     	4Bz  (3	4 r/   c                 :    | j                  || j                        S rK   )r   r   )r   r   s     r-   return_types_irzASTFunction.return_types_ir   s    $$Wdnn==r/   c                      fd}t        t         j                  |            }|D cg c]  }t         j                  |       }} j	                  ||      } j                  |      }|j                  ||      S c c}w )Nc                 ,    | j                   vxr |d uS rK   r@   pathr1   r   s     r-   <lambda>z'ASTFunction.serialize.<locals>.<lambda>       T^^!;!M r/   )r[   r    r   r!   r   r   get_function_ty)r   r   is_val	val_pathsr   r   arg_types_irret_types_irs   `       r-   	serializezASTFunction.serialize   sz     Nt~~v>?	IRS&t~~t<S	S,,Wi@++G4&&|\BB Ts   A>c                 p    fd  j                         } fd}t        t         j                   |            }d}t        |j	                               D cg c]  }|j                  |       }}|D ]s  }t         j                   |      }	 j                  j                  |g       }
|
D ]  \  }}|j                  |||        |	j                  ||      \  }}t        |||       u  j                  }|j                         D ]%  \  }}t        ||t        j                  |             ' |S c c}w )Nc                     t        | t        t        t        j                  f      r*t        j                  | D cg c]
  } |       c}|       S t        j
                  d       S c c}w rK   )rL   r[   r   r   rr   r   )rB   xmake_templates     r-   r  z.ASTFunction.deserialize.<locals>.make_template  sR    "tUH,?,?@A~~&DA}Q'7&DbII%%d++ 'Es   A&c                 ,    | j                   vxr |d uS rK   r   r   s     r-   r   z)ASTFunction.deserialize.<locals>.<lambda>  r   r/   r   )r   r[   r    rangeget_num_argsrf   r!   r   getset_arg_attrr   r"   r@   r   r   r   )r   re   ru   r   r   r   rD   r{   r   rB   
attr_specs	attr_nameattr_valrn   r@   r  s   `              @r-   deserializezASTFunction.deserialize  s"   	,
 T^^,Mt~~v>?	',R__->'?@!2771:@@ 	/D"4>>48Bb1J'1 =#	8	8<= **7F;KCdD#.	/ NN	"* 	CID#dD(*<*<S*AB	C As   D3N)rb   r   r   r   r   r   r   r   r+   r   r   r   r  r   r/   r-   r   r      sj    

 4	? tTVT[T[} >rzz >d277m >C Cr/   r   T)frozenc                   "    e Zd ZU eed<   eed<   y)BoundJITMethod__self____func__N)rb   r   r   r   __annotations__r   r   r/   r-   r  r  !  s    r/   r  c                   x   e Zd ZU dddddddddedee   dee   fdZe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j0                  ff       d Zd Zd Zej:                  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e*jV                  fdZ,d Z-defdZ.d  Z/d! Z0d" Z1d# Z2d$ Z3d% Z4d& Z5d' Z6d( Z7d) Z8d* Z9d+ Z:d, Z;e*jx                  d-e*jz                  d.e*j|                  d/e*j~                  d0e*j                  d1e*j                  d2e*j                  d3e*j                  d4e*j                  d5e*j                  d6e*j                  d7e*j                  d8iZHeeIe*j                     ef   ed9<   d: ZKd; ZLd< ZMd= ZNd> ZOd? ZPd@ ZQdA ZRe*j                  dBe*j                  dCe*j                  dDe*j                  dEe*j                  dFe*j                  dGiZYeeIe*j                     ef   edH<   dI Z[e*j                  dJe*j                  dKe*j                  dLe*j                  dMiZ`eeIe*j                     ef   edN<   dO ZbdP ZcdQ ZddR ZedS ZfdT ZgdU ZhdV ZidW ZjdX Zkdeleef   fdYZmdefdZZndpd[efd\Zod] Zpd^ Zqd_ Zrd` Zsde*j                  fdaZue*j                  dbe*j                  dciZxeeIe*j                     ef   edd<   de Zzdf Z{dg Z|dh Z}fdiZ~dj Zde*j                   ddfdkZdl ZdmdnlmZ ej*                  j
                  eej*                  j                   ee      ej
                  eej                   ee      e ee      e ee      iZeeee*j                   gef   f   edo<   xZS c c}} w )qCodeGeneratorNFr   )module	is_kernelfunction_typesr_   rA   	file_name
begin_linejit_fnr  r  c                   || _         |	| _        |	r8ddlm} t	        j
                  |      | _         || j                        | _        n7ddlm	} t        j                  |      | _         || j                        | _        d | _        || _        |dz
  | _        | j                  j                  ||d       || j                  _        || j                  _        |i n|| j                  _        |
| j                  j%                         n|
| _        |i n|| _        || _        i | _        |j/                         D ]  \  }}t1        |t2              r*|j5                  |j6                  |      | j,                  |<   @t9        |dd      }||v r't9        ||   |j6                        | j,                  |<   x|| j,                  |<    i | _        || _        |r#||j?                  d      dz   d  }tA        |d      }|| _!        || _"        d | _#        || _$        || _%        g | _&        d | _'        i | _(        | jS                         | _*        d | _+        d	| _,        y )
Nr   )GluonSemantic)TritonSemanticr#   r    r3   functionF)-contextis_gluon,triton.experimental.gluon.language._semanticr  r   GluonOpBuilderr   semantictriton.language.semanticr  r   name_loc_as_prefixr  r  set_locoptionscodegen_fns
module_mapcreate_moduler  function_ret_types	prototyper   r   rL   r   r  rb   r   r   r  rfindr.   function_namer  cur_noder_   rA   	scf_stackret_typer   _define_name_lookupdereference_namere   visiting_arg_default_value)r   r  r,  r   r.  r  r'  r(  r)  r   r  r  r  r_   rA   r  r  r  r  krt   module_names                         r-   r   zCodeGenerator.__init__)  s     R#227;DL)$,,7DM?::g.DL*4<<8DM"&"$q.Y
A6& $/ (2(:"
6<ndll002&(6(>"N"LLN 		#DAq!Z(!+

A!>A!!\26Kj(!(K)@!**!MA!"A		# )-*=*=c*BQ*F*GHM5mZPM*" , .06:6N6N6P +0'r/   builtin_namespaceprintminmaxc                 D    t        | j                  j                  ||      S rK   )r&   r  ra   )r   rd   messages      r-   _unsupportedzCodeGenerator._unsupportedr  s    +DKKOOT7KKr/   c                 r    t               }| j                  j                  ||      }||u ryt        |      ryy)NFT)objectr   r  rT   )r   r*   absent_markerrn   s       r-   _is_constexpr_globalz"CodeGenerator._is_constexpr_globalu  s6    kkoodM2-r/   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 rK   )r   r  )r*   absentr   s     r-   local_lookupz7CodeGenerator._define_name_lookup.<locals>.local_lookup  s    ;;??400r/   c                 b   j                   j                  | |      }t        ||u | j                  v t	        |      t
        u t        |t              t        |dd      t        |dd      t        |dd      j                  d      t        |dd      j                  d      t        |t        j                        t        |      j                  |       j                  t        j                   j"                  g      r|S t%        t'        j(                  d|  d	      j+                  d
d            )N__triton_builtin__F__triton_aggregate__r   r  ztriton.languagez"triton.experimental.gluon.languagez.                Cannot access global variable a   from within @jit'ed
                function. Triton kernels can only access global variables that
                are instanstiated as constexpr (`x = triton.language.constexpr(42)`). Note that this is different from
                annotating a variable as constexpr (`x: triton.language.constexpr = 42`), which is not supported.  Alternatively, set the
                envvar TRITON_ALLOW_NON_CONSTEXPR_GLOBALS=1, but we do not
                promise to support this forever.
 )r   r  r   r7  r+   r   rL   r   r   
startswithr   rS   ro   rA  r4  r   compilationallow_non_constexpr_globals	NameErrortextwrapdedentr>   )r*   rD  rn   r   s      r-   global_lookupz8CodeGenerator._define_name_lookup.<locals>.global_lookup  s   ++//$/C 6MD222I+sK0C!5u=C!7?Cr2==>OPCr2==>bcsHNN3"3'--d3 33%%AA! $ 
 %//3f 51!4 5 6=WT35GI Ir/   rI   c                     }j                   j                  fD ]  } || |      }||us|c S  t        |  d      )Nz is not defined)r7  r  rN  )r*   rD  lookup_functionrs   r@  rQ  rE  r   s       r-   name_lookupz6CodeGenerator._define_name_lookup.<locals>.name_lookup  sW    "F#/@V@V@Z@Z#Z !'f5& L! tfO455r/   )strr?  r   )r   rT  r@  rQ  rE  s   ` @@@r-   r2  z!CodeGenerator._define_name_lookup  sG    	1s 	1	I 	IB 	6c 	6c 	6 	6 r/   c              #   0   K   || _         d  d | _         y wrK   )r%  )r   prefixs     r-   _name_loc_prefixzCodeGenerator._name_loc_prefix  s     "("&s   c                    t        |t        j                  t        j                  f      r:|j	                  | j
                  j                  ||j                                      y t        |      rTg }|j                  |       |D ];  }|j	                  | j
                  j                  ||j                                      = y y rK   )
rL   r   rs   block_argumentr&  r   create_name_locget_locrN   rz   )r   rn   r*   r{   handles        r-   _maybe_set_loc_to_namez$CodeGenerator._maybe_set_loc_to_name  s    cBHHb&7&789KK44T3;;=IJc"GOOG$! Ut||;;D&..BRSTU #r/   r*   rs   rI   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   )r   r*   rs   s      r-   	set_valuezCodeGenerator.set_value  s      "D %r/   c                 r    | j                   j                         }| j                   j                         }||fS rK   )r   r\  r   )r   locips      r-   _get_insertion_point_and_locz*CodeGenerator._get_insertion_point_and_loc  s1     ll""$\\--/3wr/   c                 p    | j                   j                  |       | j                   j                  |       y rK   )r   r   r&  )r   rc  rb  s      r-   _set_insertion_point_and_locz*CodeGenerator._set_insertion_point_and_loc  s&    ,,R0S!r/   c                    | j                   j                         }| j                   j                  |       | j                  j	                  |       | j                  |j                         | j                  j                          |j                          g }g }g }|j                         D ]  \  }}t        |      r~| j                  |   }	| j                  ||	|       t        |g      }
t        |	g      }|
|k7  sQ|j	                  |       |j	                  |j                         |j                  |
       || j                   vrJ d| d        |j#                         | _        i | _        |||fS )NLoop carried variable z is not a triton value)r   create_blockset_insertion_point_to_startr0  appendvisit_compound_statementr   poperaser   rN   r   _verify_loop_carried_variabler|   r+   extendr   copy)r   rd   r   blockinit_tysinit_handlesnamesr*   live_valloop_vallive_handlesloop_handless               r-   _find_carrieszCodeGenerator._find_carries  sF   ))+11%8d#%%dii0
 %mmo 	jND();;t,2248L3XJ?3XJ?</LL&OOHMM2 ''54??2i6LTFRh4ii2	j llnlH,,r/   c                     t        |      s|g}|D ]/  }| j                  |       t        |t        j                        s/ y  y rK   )r\   r   rL   r   r   )r   stmtsstmts      r-   rl  z&CodeGenerator.visit_compound_statement  s?    U#GE 	DJJt $

+	r/   c                 D    t         j                  j                  | |       y rK   r   NodeVisitorr   r   s     r-   r   zCodeGenerator.visit_Module
      %%dD1r/   c                     | j                  |j                        }|J t        j                  |j                  D cg c]  }| j                  |       c}      }|S c c}w rK   )r   r   r   r   elts)r   rd   r   eltr  s        r-   
visit_ListzCodeGenerator.visit_List  sN    jj"{{~~$))D3tzz#DE Es   Ard   c                    t        |j                        dk7  rt        d      |j                  d   }| j                  |j                        }t        |t              st        d      g }|D ]R  }| j                  |j                  j                  |       |j                  | j                  |j                               T t        |      S )Nr#   z'nested comprehensions are not supportedr   z'only tuple comprehensions are supported)r   
generators
ValueErrorr   iterrL   tl_tupleNotImplementedErrorr`  targetr   rk  r  )r   rd   compr  resultsr   s         r-   visit_ListCompzCodeGenerator.visit_ListComp  s    t1$FGGq!zz$))$$)%&OPP 	1DNN4;;>>40NN4::dhh/0	1   r/   c                 &     j                  |j                        }g } fd |      }|t        j                  }nCt	        |t        j
                  j                        sJ |j                  |       |j                  } j                  j                  |        j                  | _        n* j                  |k7  rt        d j                   d|        j                  j                         } j                  j                  |       y )Nc                     t        | t        j                        rt        |       S t        | t        j                  t
        t        f      rj                  j                  |       S | S rK   )	rL   r   r   rw   r   intfloatr#  	to_tensor)rs   decayr   s    r-   r  z)CodeGenerator.visit_Return.<locals>.decay'  sO    %0-eU;;EH$6$6U#CD}}..u55Lr/   zInconsistent return types:  and )r   rs   r   voidrL   rR   r   rz   r+   r   rF   r1  	TypeErrorri  set_insertion_point_to_end)r   rd   	ret_valuer{   ret_typost_ret_blockr  s   `     @r-   r   zCodeGenerator.visit_Return#  s    JJtzz*		 )$	]]Fi)A)ABBB!!'*^^F!== "DM]]f$9$--fXVWW 224//?r/   c                     | j                  |j                        }t        |t        j                  j
                        sJ |j                  S rK   )r   rs   rL   r   rR   r   rx   )r   rd   rf   s      r-   visit_StarredzCodeGenerator.visit_StarredA  s8    zz$**%$ 3 3444{{r/   c                 @	   | j                  |j                        \  }}| j                  r| j                  |d      t	        |j                  j
                  d d d         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/                         }| j                  j1                  | j                        }| j2                  0| j2                  j5                  | j                  | j"                         t7        ||      D ])  \  }}| j9                  ||       | j;                  ||       + | j"                  j=                         }| j"                  j?                  |       | jA                  |jB                         | j"                  j=                         jE                         rJ | jF                  | jF                  tH        jJ                  k(  r2tH        jJ                  | _#        | j"                  jM                  g        n	tO        | jF                  tH        jP                        r&| jF                  jR                  | j                  _*        n| jF                  g| j                  _*        | j                  jW                  | j                  j!                  | j"                               | j"                  jM                  | j                  jY                  | j"                        D cg c]  }| j"                  j[                  |       c}       | j                  j]                          |r| j"                  j_                  |       y y # d| _        w xY wc c}w )Nz,nested function definition is not supported.r#   )r   r   targetsrs   )r  rs   
annotationTFpublicprivate)0r   rf   re   r=  r`   defaultsr  rh   r   r   r   r   	AnnAssignr4  r  r,  r   r   get_or_insert_functionr  r.  r_   	push_backadd_entry_blockr  rA   initialize_calleezipr^  r`  r   rj  rl  r   has_terminatorr1  r   r  rF   rL   rr   rv   r   
reset_typer   create_poisonfinalizer  )r   rd   rc   kwarg_namesrD   default_valuearg_noder  r*   	st_target	init_node
visibilityfn_tyentry
arg_valuesarg_name	arg_value	insert_ptrB   s                      r-   r   zCodeGenerator.visit_FunctionDefF  sf   !%DII!6	;77##D*XYY )$))*<*<TrT*B C 	8A}yy~~qb1f-H!,,J<<DDciik:I!JJ	{-P	MM-\fg	8::::26/

9%27/	8" "&XY
((6,,55dkk4CUCUW\^hjnjwjwxdgg&'')^^//8
*11$''4<<H#&y*#= 	0Hi''	8<NN8Y/	0 LL446	11%8%%dii0 <<335DDFFF== DMMX]]$B$MMDMLLR $--)<)<=+/==+>+>(,0MM?(GGt~~77EFLLt~~GeGefjfrfrGstdll88<tuLL33I> E 38/> us   :&R-"R	Rc                     g }|j                   D ]  }|| j                  |      gz  } | j                  |j                        }||fS rK   )rf   r   kwarg)r   rd   rc   rh   r  s        r-   visit_argumentszCodeGenerator.visit_arguments~  sL    	99 	+C$**S/**I	+jj,+%%r/   c                 Z    t         j                  j                  | |       |j                  S rK   )r   r  r   rh   r   s     r-   	visit_argzCodeGenerator.visit_arg  s    %%dD1xxr/   c                 b   | j                  |j                        }| j                  |j                        }| j                  |j                        }|t        k(  rE|| j
                  v rt        | d      t	        |      }|| j
                  |<   | j
                  |   S | j                  |      S )Nz4 is already defined. constexpr cannot be reassigned.)r   r  r  rs   r   r   r  r   )r   rd   r  r  rs   s        r-   visit_AnnAssignzCodeGenerator.visit_AnnAssign  s    ZZ0
DKK(

4::&"$ F8 ,D "E F Fe$E"'DKK;;v&&  &&r/   c                    t        |j                  t        j                        sJ t        |t        j                        r| j                  ||      S t        |t        j                        r=t        |j                        D ]$  \  }}| j                  ||j                  |          & y t        |t        j                        rt        d      t        |t        j                        sJ | j                  | j                  |      |       y )Nz/Attribute assignment is not supported in triton)rL   r   r   r   	Subscriptvisit_Subscript_Storer	   r`   r  assignTargetrx   r   r  r   r`  r   )r   r  rs   rD   s       r-   r  zCodeGenerator.assignTarget  s    &**cii000fcmm,--fe<<fcii(&v{{3 ;	6!!&%,,q/:;fcmm,%&WXX&#((+++tzz&)51r/   c                      fdt        |t        j                        r|j                  gn|j                  }t        |      dk(  sJ |d   }t        |t        j                        rF j                  |j                        5    j                  |j                              }d d d        n!  j                  |j                              } j                  |       y # 1 sw Y   xY w)Nc                    t        | t        j                        rt        |       S t        j                  t        j                  f}t        |       } | 2t        |       s't        | |      sj                  j                  |       } | S rK   )	rL   r   r   rw   rS   r   rN   r#  r  )rs   native_nontensor_types_sanitize_valuer   s     r-   r  z3CodeGenerator.visit_Assign.<locals>._sanitize_value  sm    %0-e_EE&.nnhnn%E"(/E $U+u&<=//6Lr/   r#   r   )rL   r   r  r  r  r   r   rX  r   r   rs   r  )r   rd   r  r  rx   r  s   `    @r-   r   zCodeGenerator.visit_Assign  s    		 $.dCMM#B4;;-7|q   fchh'&&vyy1 A(DJJ)?@A A %TZZ

%;<F&&)	A As   "C%%C.c                 P   t        j                  |j                        }t        j                         |_        t        j                  ||j                  |j                        }t        j                  |j                  g|      }| j                  |       | j                  |      S )Nr  )rq  deepcopyr  r   Loadr   BinOpoprs   r   r   )r   rd   lhsrhsassigns        r-   r   zCodeGenerator.visit_AugAssign  sk    mmDKK(((*iiTWWdjj1T[[M=

6zz#r/   c                     t        |j                        t        j                  u r|j                  S | j                  |j                        S rK   )r+   r   r   r   r   r3  r   s     r-   r   zCodeGenerator.visit_Name  s4    >SYY&77N$$TWW--r/   c                 D    t         j                  j                  | |       y rK   r  r   s     r-   visit_StorezCodeGenerator.visit_Store  r  r/   c                 D    t         j                  j                  | |       y rK   r  r   s     r-   
visit_LoadzCodeGenerator.visit_Load  r  r/   c                     |j                   D cg c]  }| j                  |       }}t        j                  |      S c c}w rK   )r  r   r   r   )r   rd   r  rf   s       r-   visit_TuplezCodeGenerator.visit_Tuple  s5    '+yy1!

111~~d## 2s   >c                 n   t        |      r t        ||      || j                        S t        |      r5t        j                  dd|      } t        ||      || j                        S t        |t        t        j                  f      st        |t              rt        |      } t        ||      |      S )N	_semanticz__(.*)__z__r\1__)	rP   r   r#  r(   subrL   r   r   r   )r   method_namer  r  reverse_method_names        r-   _apply_binary_methodz"CodeGenerator._apply_binary_method  s    S!,73,SDMMJJS!"$&&j+"N473 34SDMMRR#	8>>:;
3PY@ZC.C(wsK(--r/   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  r+   r  r=  formatrb   r  )r   rd   r  r  r  s        r-   visit_BinOpzCodeGenerator.visit_BinOp  s    jj#jj$2266tDGG}E##D$^$e$efjfmfmfvfv$wy y((c3??r/   __add____sub____mul____truediv____floordiv____mod____pow__
__lshift__
__rshift____and____or____xor__r  c                 R   | j                   j                  |       | j                  |j                         | j                   j	                         }| j
                  j                         }| j                  j                         }i }|j                         }|j                  r| j                   j                  |       |j                         | _        i | _        | j                  |j                         | j
                  j                         }| j                   j	                         }| j                  j                         }g }	|j                         D ]  \  }
}t        |      st        ||
   g      }t        ||
   g      }||k(  r6|	j                  |
       ||
   ||
<   ||
   ||
<   |df|dffD ]S  \  }}t        ||
         t        |      k(  }|r||
   j                  |j                  k(  r?J d|
 d| d| d||
             t        |j                         |j                         z        D ]i  }
|
|	v r||
   }|j                  }||
   }|j                  }t        |      t        |      k(  }|r||k(  sJ d|
 d| d	| d
       |	j                  |
       k |||||	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   rj  rl  r   r   r   rq  r   r   r   rN   r|   rk  r+   r<   keys)r   rd   r   
then_block
else_block	then_defs	then_vals	else_defs	else_valsru  r*   rs   then_handleselse_handlesdefs
block_name
type_equalthen_valthen_tyelse_valelse_tys                        r-   visit_then_else_blocksz$CodeGenerator.visit_then_else_blocks  s   11*=%%dii0\\557
OO((*	KK$$&		LLN	;;LL55jA!,,.DK DO))$++6,,.I99;J((*I "==? 	OKD%#E*/40ABL/40ABL|+LL'oIdO'oIdO&/%89f:M$N O j!$t*-e<
!d4joo&C O)$}UG D)l*A$t*OOCO	O( 9>>+inn.>>? 	Du} HmmG HmmGh4>9J'W"4 .&tf,A' K##*)1..4 LL	 )ZUBBr/   c                 @   t        |       5 }|\  }}| j                  j                         }| j                  j                         }| j                  j                  |       | j                  j	                  |j
                  ||       | j                  ||||      \  }}}| j                  j                         }	| j                  j                  |       |j                         rJ |        t        fd|D              }
| j                  j                  |	|
       | j                  j                  |       |j                         rJ |        t        fd|D              }| j                  j                  |	|       t        |
      t        |      k(  sJ t        |
|      D ];  \  }}|j                         }||j                         k(  sJ |	j                  |       = 	 d d d        | j                  j                  	       t        t        
            D cg c]  }|	j!                  |       }}D cg c]  }|   j"                   }}t%        ||      }t        ||      D ]  \  }}| j'                  ||        y # 1 sw Y   xY wc c}w c c}w )Nc              3   (   K   | ]	  }|     y wrK   r   r   r*   r  s     r-   r   z3CodeGenerator.visit_if_top_level.<locals>.<genexpr>@       /RD	$/R   c              3   (   K   | ]	  }|     y wrK   r   r   r*   r  s     r-   r   z3CodeGenerator.visit_if_top_level.<locals>.<genexpr>E  r  r  )r   r   ri  r  create_cond_branchr]  r  r  r|   create_branchr   r  get_typeadd_argumentrj  r  rh   r+   r   r`  )r   condrd   srr   ip_blockr  r  ru  endif_blockr  r   then_helse_hrB   rD   res_handlesr*   rv   
new_values	new_valuer  r  s                        @@r-   visit_if_top_levelz CodeGenerator.visit_if_top_level0  sI   d# 	-r "GX224J224JLL33H=LL++DKKZP ++D':zR @Iy*j% ,,335KLL33J?!002CzlC2//RE/RRLLL&&{LALL33J?!002CzlC2//RE/RRLLL&&{LA|$L(9999"%lL"A -__&V__....((,-/	-: 	11+>38\9J3KLa{q)LL278$4%%88(e<
"5*5 	,OD)NN4+	,E	- 	-> M8s   GJ
!J?J
Jc                    t        |       5 }|\  }}| j                         \  }}| j                  j                         }|j                  r| j                  j                         nd }	| j                  ||||	      \  }}	}
t        fd|
D              }t        |
|      D ]  \  }}| j                  ||        | j                  ||       | j                  j                  |D cg c]  }|j                          c}|j                  d      }|j                  |j                                | j                  j                  |j                                t!        |
      dkD  r| j                  j#                  |       |j                  s|j%                         }	n|	j                  |j%                                | j                  j                  |j%                                t!        |
      dkD  rUt        fd|
D              }t        |
|      D ]  \  }}| j                  ||        | j                  j#                  |       d d d        t'        t!                    D cg c]  }j)                  |       }}
D cg c]  }|   j*                   }}t-        ||      }t        |
|      D ]  \  }}| j/                  ||        y c c}w # 1 sw Y   xY wc c}w c c}w )Nc              3   (   K   | ]	  }|     y wrK   r   r  s     r-   r   z-CodeGenerator.visit_if_scf.<locals>.<genexpr>`  r  r  Tr   c              3   (   K   | ]	  }|     y wrK   r   r  s     r-   r   z-CodeGenerator.visit_if_scf.<locals>.<genexpr>o  s     3VIdO3Vr  )r   rd  r   ri  r   r  r|   r  r^  rf  create_if_opr  r]  merge_block_beforeget_then_blockr  r   create_yield_opget_else_blockr  
get_resultr+   r   r`  )r   r  rd   r  r   r1   rc  last_locr  r  ru  r  r*   rn   hif_opr   rD   r  rv   r  r  r  r  s                         @@r-   visit_if_scfzCodeGenerator.visit_if_scfW  s   d# 	;rJGQ<<>LB224J8<224J++D':zR @Iy*j% 0/RE/RRL 5 7	c++C67--b(;LL--\.Rqzz|.RTXT_T_aefE))%*>*>*@ALL33E4H4H4JK5zA~,,\:;;"113
--e.B.B.DELL33E4H4H4JK5zA~33VPU3VV!$UL!9 ;ID#//T:;,,\:5	;8 5:#l:K4LMqu''*MM278$4%%88(e<
"5*5 	,OD)NN4+	,' /S	; 	;8 N8s+   CK(K
?EK)K#K(KK c           	         | j                  |j                        }t        |      r+t        |      r| j	                  |d      |j
                  j                         rat        j                  dt        j                  |j                        z         t        j                  j                  || j                  |       }|j                  t        j                   | j                        }t#        | j$                        j                  |      r1| j&                  r| j	                  |d      | j)                  ||       y | j+                  ||       y t-        |      }t        |      t.        vrO| j	                  |dj1                  dj3                  d t.        D              t        |      j4                              |r|j6                  n|j8                  }| j;                  |       y )	Nz=Boolean value of Tensor with more than one value is ambiguousziIf conditional called with multidimensional Tensor instead of scalar; please use "if (%s).item()" instead)r  
_generatorr  zMCannot have `return` statements inside `while` or `for` statements in triton.O`if` conditionals can only accept values of type {{{}}}, not objects of type {}, c              3   4   K   | ]  }|j                     y wrK   rb   r   r1   s     r-   r   z)CodeGenerator.visit_If.<locals>.<genexpr>       !G!**!G   )r   testrP   rY   r=  r+   rW   warningswarnr   unparser   rR   _unsplatr#  toint1r   r   r0  r  r*  r   _condition_typesr  r:   rb   r   r   rl  )r   rd   r  active_blocks       r-   r   zCodeGenerator.visit_Ifz  s   zz$))$T"$T*''.mnnyy!!# Bkk$)),-.  }}--ddmmX\-]778==DMM7BD$T[[177=>>++mo o''d3!!$-'-DDz!11''krr		!G6F!GGT
++-. .
 )-499$++L)),7r/   c           	      8   | j                  |j                        }t        |      rE|j                  t        j
                  | j                        }t        |       5  | j                         \  }}| j                  j                         }| j                  j                  |       | j                  j                  | j                  |j                              }| j                  j                         }| j                  j                         }| j                  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        j6                  j9                  |j;                  d      |	      nd cd d d        S t=        |      }t#        |      t>        vrO| jA                  |djC                  djE                  d t>        D              t#        |      jF                              |r| j                  |j                        S | j                  |j                        S # 1 sw Y   y xY w)	Nr  zATernary expression with dynamic condition has inconsistent types r  Tr   r-  r.  c              3   4   K   | ]  }|j                     y wrK   r0  r1  s     r-   r   z,CodeGenerator.visit_IfExp.<locals>.<genexpr>  r2  r3  )$r   r4  rP   r9  r   r:  r#  r   rd  r   ri  rj  r  r   r   r   rf  r+   r  to_irr!  r]  r"  r#  r  r$  r%  rR   r   r&  r   r;  r=  r  r:   rb   )r   rd   r  rc  r'  r  r  r  r  r1  ret_type_irr)  s               r-   r   zCodeGenerator.visit_IfExp  s   zz$))$T"778==DMM7BD!$' !d#@@BH!\\668
99*E==224::dii3HI!\\==?
!\\668
99*E  ==224::dkk3JK!\\==?
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Pc                     t        |j                        dk(  r|j                  d   j                  } j                  |j                        }|t
        j                  k(  r_|j                  D cg c]  } j                  |       }} ||d j                  i5   j                  |j                         d d d        y g }|j                  D ]  }|j                  } j                  |j                        }	|j                  D cg c]  } j                  |       }}t         fd|j                  D              }
 |	|d j                  i|
}|j                  |        t        ||j                        D ]O  \  }}|j!                         }|j"                  # j                  |j"                        } j%                  ||       Q t'         j(                        j                  |      r j+                  |d       j                  |j                         t-        |      D ]  }|j/                  d d d         y c c}w # 1 sw Y   y xY wc c}w )Nr#   r   _builderc              3   @   K   | ]  }j                  |        y wrK   r   )r   kwr   s     r-   r   z+CodeGenerator.visit_With.<locals>.<genexpr>  s     >"tzz"~>r   r  zCCannot have `return` statements inside `with` statements in triton )r   r   context_exprr   r   r   
async_taskrf   r   rl  r   dictkeywordsr#  rk  r  r   optional_varsr`  r   r   r=  reversedr   )r   rd   r  withitemClassrh   rf   cm_listr   callre   kwscmresvar_names   `             r-   
visit_WithzCodeGenerator.visit_With  s    tzz?ajjm00G JJw||4M 3 333:<<@C

3@@"D@4<<@ =11$))<=JJ 	D$$DDII&B/3yy9DJJsO9D9>>>CT:T]]:c:BNN2	 GTZZ0 	.HB,,.C!!-::d&8&89x-		.
 !-33D9##D*opp%%dii07# 	*BKKdD)	*+ A= :s   /IIIIc                      y rK   r   r   s     r-   
visit_PasszCodeGenerator.visit_Pass  s    r/   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                  u rt        ||u       S t        |j                  d         t        j                  u 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.)r   comparatorsopsr=  r   r  r   r+   r   Isr   IsNot_method_name_for_comp_opr  r  rb   r  )r   rd   r  r  	lhs_value	rhs_valuer  s          r-   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??r/   __eq____ne____lt____le____gt____ge__rZ  c           
         | j                  |j                        }| j                  j                  t	        |j
                              }|*| j                  |d|j
                  j                   d      t        |      r t        ||      | j                        S 	  t        ||             S # t        $ r? |dk(  rt        |       cY S | j                  |d| dt	        |      j                         w xY w)NzAST unary operator 'z!' is not (currently) implemented.r  __not__z)' is not (currently) implemented on type )r   operand_method_name_for_unary_opr  r+   r  r=  rb   rP   r   r#  AttributeErrorr   )r   rd   rf  re   s       r-   visit_UnaryOpzCodeGenerator.visit_UnaryOp  s    **T\\*++//TWW>:##D,@AQAQ@RRs*tuuW%'77B'$--@@	t'77B')) 	tY W--##,RD0YZ^_fZgZpZpYqrt t	ts   B0 0C8,C8__neg____pos__re  
__invert__rg  c           	      `   t        |      sJ d| d       t        |      sJ d| d       t        |      t        |      u s"J d| dt        |       dt        |              t        |      r?|j                  |j                  k(  s%J d| d|j                   d|j                   d	       y y )
Nzcannot reassign constexpr z in the looprh  z changed type, was z but is now zLoop-carried variable z has initial type z but is re-assigned to z: in loop! Please make sure that the type stays consistent.)rN   r+   rP   )r   r*   rw  rv  s       r-   ro  z+CodeGenerator._verify_loop_carried_variable  s    )Z-Gv\+ZZ))Z-Gv\+ZZ)H~h/ 	l$TF*=d8n=M\Z^_gZhYij	l/$X.(--8==2P 	@$TF*<X]]O L%%-]]O 4?@	@P2P.r/   c                 8    | j                  |j                        S rK   )r   rE  r   s     r-   visit_withitemzCodeGenerator.visit_withitem"  s    zz$++,,r/   c                     t               5 }|\  }} j                         \  }} j                  ||      \  }}}	|D 
cg c]  }
|
j                          }}
 j	                  ||        j
                  j                  ||      } j
                  j                  |j                         |      } j
                  j                  |       t        t        |            D cg c]  }|j                  |       }}t        ||	      }t        ||      D ]5  \  }}| j                  |<   | j                   |<    j#                  ||       7  j%                  |j&                        }t)        |t*        j,                        rB|j.                  r*|j1                  d j
                  j3                                |j,                  } j
                  j5                  |        j
                  j7                  |j8                  |        j
                  j                  |j;                         |      } j
                  j                  |       t        t        |            D cg c]  }|j                  |       }}t        ||	      }t        ||      D ]5  \  }}| j                  |<   | j                   |<    j#                  ||       7  j<                  j?                  |        jA                  |jB                          j<                  jE                          tG         fd|D              } j
                  jI                  |       d d d        t        t                    D cg c]  }jK                  |       }}t        |	      }t        |      D ]5  \  }}| j                  |<   | j                   |<    j#                  ||       7 |jL                  D ]  }J d        y c c}
w c c}w c c}w # 1 sw Y   xY wc c}w )Nllvm.loop_annotationc              3   <   K   | ]  }j                   |     y wrK   r   r   r*   r   s     r-   r   z,CodeGenerator.visit_While.<locals>.<genexpr>O       0UtT1B0U   zNot implemented)*r   rd  rz  r  rf  r   create_while_opcreate_block_with_parent
get_beforerj  r  r   rh   r   r  r   r   r^  r   r4  rL   r   	conditiondisable_licmset_attrget_disable_loop_licm_attrr  create_condition_opr]  	get_afterr0  rk  rl  r   rm  r|   r$  r&  r   r   r  r   )r   rd   r  r   r   rc  r'  ru  rt  init_fe_tysr(  rs  while_opbefore_blockrD   
block_argscondition_argsr*   rn   r  after_blockbody_handles	body_argsyield_handlesresult_handlesresult_valsnew_defr}  s   `                           r-   visit_WhilezCodeGenerator.visit_While%  s{   d# *	8r$&!G\<<>LB/3/A/A$/P,E<.:;

;H;--b(;||33HlKH<<@@ATATAVX`aLLL55lC7<S=N7OP!,**1-PJP0[IN 7 7	c$'D!(+%++C67 ::dii(D$ 2 23$$%%&<dll>e>e>gh~~LL33LALL,,T[[*E,,??@R@R@TV^_K LL55kB8=c,>O8PQ1KOOA.QLQ+L+FI 	2 7	c$'D!(+%++C67 NN!!$'))$))4NN 00Uu0UUMLL((7U*	8Z ;@L@Q:RSQ(--a0SS).+F 4 	7MD' 'DKK$+DOOD!''6	7
 KK 	6D+++5	6] < Q& R?*	8 *	8Z Ts>   3O7 O(BO7"O-:E%O7O27CO7&P(O77P c                    t        |j                  t        j                        sJ | j	                  |j
                        }| j	                  |j                        }t        |      r | j                  ||j                  ||gi       S ||   S rK   )
rL   r   r   r  r   rs   slicerN   call_Method__getitem__)r   rd   r  slicess       r-   visit_Subscript_Loadz"CodeGenerator.visit_Subscript_Load^  so    $((CHH---jj$DJJ'C ##D#//3"MM6{r/   c                     t        d      )Nz&__setitem__ is not supported in triton)r  )r   rd   rs   s      r-   r  z#CodeGenerator.visit_Subscript_Storef  s    !"JKKr/   c                 $    | j                  |      S rK   )r  r   s     r-   visit_SubscriptzCodeGenerator.visit_Subscripti  s    ((..r/   c                 ^    |j                   D cg c]  }| j                  |       c}S c c}w rK   )dimsr   )r   rd   dims      r-   visit_ExtSlicezCodeGenerator.visit_ExtSlicel  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 }
d }d}d}d}d}|t        j                  u ru ||i |}|j                  }|j                  }|j                  }|j.                  }
|j0                  }|j2                  }|j4                  }|j6                  }|j8                  }n|t        u rt;        |      dkD  r|d   n# j                  t)        j<                  d            }t;        |      dkD  r|d   n' j                  |j                  j                  d         }t;        |      dkD  r|d   n# j                  t)        j<                  d            }nt?        d      d}tA        |      r+|j                  dk  rt        |j                         }d}||}} jB                  jE                  |      } jB                  jE                  |      } jB                  jE                  |      }|jF                  jI                         r4|jF                  jI                         r|jF                  jI                         s3tK        d|jF                   d	|jF                   d	|jF                   d
       jB                  jM                  |jF                  |jF                        } jB                  jM                  ||jF                        }|jO                   jP                        }|jR                  t        jT                  jF                  jV                  jX                  k(  }|jZ                  }|jZ                  }|jZ                  } jP                  j]                  |||      } jP                  j]                  |||      } jP                  j]                  |||      } jP                  j_                  |      } ja                  |j                  j                   t        jT                  jc                  ||             te               5 }|\  }} jg                         \  }} ji                  ||      \  }}} jk                  ||        jP                  jm                  ||||      } to        |
      +| jq                  d jP                  js                  |
             to        |      +| jq                  d jP                  js                  |             |r*| jq                  d jP                  ju                                |r*| jq                  d jP                  ju                                |r*| jq                  d jP                  ju                                |r*| jq                  d jP                  jw                                 jx                  j{                  |       | j}                  d      }! jP                  j                  |!       t        t;        |            D cg c]  }|!j                  |dz          }"}t        |"|      }#t        ||#      D ])  \  }$}% j                  |%|$        ja                  |$|%       +  j#                  |j$                          jx                  j                          t         fd|D              }&t;        |&      dkD  r jP                  j                  |&       |!j                         }'|'j                         dk(  sJ d        jP                  j                  |!       | j                         }|r8 jP                  j                  ||      } jP                  j                  ||      } j                  |j                  j                      jZ                  j                  |        ja                  |j                  j                   t        jT                  jc                  ||              j                  ||j                  j                          d d d        t        t;                    D cg c]  } j                  |       }(}t        |(      })t        |)      D ])  \  }$}% ja                  |$|%        j                  |%|$       + |j&                  D ]  }	J d        y c c}w c c}w # 1 sw Y   xY wc c}w )Nc              3   @   K   | ]  }j                  |        y wrK   r   r   keywordr   s     r-   r   z*CodeGenerator.visit_For.<locals>.<genexpr>r  s     Q74::g.Qr   Fr#   r   r   zAOnly `range` and `static_range` iterators are currently supportedTz0For loop bounds and step must all be ints, are (r.  r  ztt.num_stagesztt.loop_unroll_factorztt.disallow_acc_multi_bufferz
tt.flattenztt.warp_specializerq  c              3   <   K   | ]  }j                   |     y wrK   rs  rt  s     r-   r   z*CodeGenerator.visit_For.<locals>.<genexpr>  ru  rv  z7We use SCF, so the loop body should only have one blockz)Don't know what to do with else after for)Nr   r  r   rf   rG  rH  r   static_ranger  startrs   endstepr   r   r  r   rl  r   r   r   r  r   
num_stagesloop_unroll_factordisallow_acc_multi_bufferflattenwarp_specializer{  r   NumRuntimeErrorrT   r#  r  rS   is_intr  integer_promote_implr?  r   int_signednessrR   
SIGNEDNESSSIGNEDr]  create_int_castr  r`  r   r   rd  rz  rf  create_for_opr   r|  get_int32_attrget_unit_attrr}  r0  rk  get_bodyrj  rh   r   r  r^  rm  r|   r$  
get_parentsizeget_induction_var
create_sub
create_addreplace_all_uses_withr&  )*r   rd   IteratorClassrh   	iter_argsiter_kwargsiteratorr  rD   r}  r  r  r  r  r  r{  lbubr  negative_stepiv_type
iv_ir_typeiv_is_signedivr  r   r   rc  r'  ru  rt  rs  for_opfor_op_bodyblock_handlesr  r*   rn   r  for_op_regionr  result_valuess*   `                                         r-   	visit_ForzCodeGenerator.visit_Foro  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!)!<!<(0(J(J%&&G&66O#00Le# "%Y!!31CGGAJ9OB!$Y!!31DIINNSTDU9VB#&y>A#59Q<4::cggaj;QDbcc4::>djj[)D MB]]$$R(]]$$R(}}&&t,xx (9ARARATNrxxjXZ[][c[cZddfgkgqgqfrrstuu--44RXXrxxH--44WdjjI]]4<<0
--1D1D1O1O1V1VVYYYY{{\\))"j,G\\))"j,G||++D*lK\\''
3t{{~~x}}';';B'HId# 0	<r$&!G\<<>LB,0,>,>tW,M)E< --b(;\\//BlKF#J/;1L1LZ1XY#$67C 79T9TUg9hi( >@Z@Z@\]dll.H.H.JK 4dll6P6P6RS 68_8_8abNN!!$' //!,KLL55kB=B3|CT=UV[__QU3VMV,]HEJ 
3 *	c++C6tS)* ))$))4NN 00Uu0UUM =!A%,,];'224M %%'1,g.gg, LL55kB))+B\\,,R4\\,,R4KK'..DDRHNN4;;>>8==+?+?G+LM''DKKNN;a0	<f 9>c,>O8PQ1&++A.QQ+NHEUM2 	3ID#NN4%''T2	3 KK 	6DEEE5	6{ @z W30	< 0	<f Rs+   gGg(g	G-gg	ggc                     | j                  |j                        }| j                  |j                        }| j                  |j                        }t	        j
                  |||      S rK   )r   lowerupperr  r   r  )r   rd   r  r  r  s        r-   visit_SlicezCodeGenerator.visit_Slice  sK    

4::&

4::&zz$))$~~eUD11r/   c                 8    | j                  |j                        S rK   )r   rs   r   s     r-   visit_IndexzCodeGenerator.visit_Index  s    zz$**%%r/   c                 P    |j                   | j                  |j                        fS rK   )rh   r   rs   r   s     r-   visit_keywordzCodeGenerator.visit_keyword  s    xxDJJ///r/   c                     | j                  |j                        }|j                  | j                  |j                        nd}t        j                  j                  ||| j                        S )Nr  r  )r   r4  msgr   rR   device_assertr#  )r   rd   r4  r  s       r-   visit_AssertzCodeGenerator.visit_Assert  sQ    zz$))$&*hh&:djj"}}**4*NNr/   re   c                    t        j                  |j                  g|i |}|j                  D cg c]  }||   	 }}t	        |      D ]W  \  }}t        |t        j                  t        t        t        t        f      s6t        j                  j                  |      ||<   Y t        |d       }|D 	ci c]  }	|	t        ||	       }}	t        |d       }
|
D 	cg c]  }	t        ||	       }}	|xs | j                   }t#        t%        |      |D cg c]  }|j&                   c}||      }| j(                  j+                  |      sZt-        |      \  }}|D cg c]Y  }|/t        |t        t        t        j                  j                  f      rt        j                  j                  n|j&                  [ }}t/        g ||t1                     }t3        | j4                  ||j7                         | j(                  ||| j8                  |j:                  ||| j<                  j>                  | j<                  j@                  | j<                  jB                  || jD                        }	 |jG                  |jI                                |jZ                  }|| j8                  |<   n| j8                  |   }| j(                  j]                  |      }t_        |      }| j<                  ja                  ||      }|t        jb                  k(  ry te        |jg                               D cg c]  }|ji                  |       }}tk        tm        ||g            S c c}w c c}	w c c}	w c c}w c c}w # tJ        $ rL}tL        jN                  jP                  r tS        | jT                  jV                  | jX                  d       |d }~ww xY wc c}w )Nc                     t        |      S rK   rT   r1   r  s     r-   r   z0CodeGenerator.call_JitFunction.<locals>.<lambda>	  s    M!4D r/   c                     t        |       S rK   r  r  s     r-   r   z0CodeGenerator.call_JitFunction.<locals>.<lambda>  s    q9I5I r/   )r  r  r.  r  r_   r  r  r'  r(  r)  rA   r   )7inspectgetcallargsre   rc   r`   rL   r   rS   r  r  r   r   rR   r   r    r!   rA   rG   r   r+   r  has_functionr   r   rG  r  r  get_capture_scoper+  r_   r   r'  r(  r)  r   r   parse	Exceptionr   rL  front_end_debuggingr$   r  ra   r/  r1  get_functionr|   rM  r  r  get_num_resultsr&  nextr   )r   re   rf   r   rA   r*   rD   rh   args_cstr   	args_pathargs_valfn_namer  r  r   r,  r   ecallee_ret_typesymbolcall_opr{   s                          r-   call_JitFunctionzCodeGenerator.call_JitFunction  sC   ""255:4:6:')||4tT
44o 	7FAs#sD+NO"--11#6Q	7 !'DEDLMDD+D$77MM!$(IJ	>GHd%dD1HH'>4+>+>M"-H/MS/MxYgh{{''0$8$<!Iz    ,/;*SFJCQYQ^Q^QdQdEe;g''loltltuI 
 $B	8TVDI%dllIr?S?S?U^b^i^irt4;DLcLc/1{{i\f.2ll.B.BPTP\P\PhPh151H1HYg/3}}>IT
+ (00O/>D##G,"55g>O))'2'1,,##FH5hmm+278O8O8Q2RSQ7%%a(SS'/1BCDD[ 5
 NH 0N
  T$$88&t{{tLRSS	T Ts=   M6MM!M&
AM+,M0 'O0	O9AO  Oc                 x   t        |t        t        f      r(|j                  d|j                         |j
                  }t        |t              r t        |||       | j                  |||      S t        |d      rt        |j                        s/t        j                  j                  |      st        |t              rt               }t        |t              r t!        j"                  |j$                        }nt!        j"                  |      }d|j&                  v r| j(                  |d<   d|j&                  v r| |d<   	  ||i ||}t        |t*              rt        j*                  |      }|S || j<                  j?                         v rtA        tB        |      } ||i |}d }	t        |tD        j*                  t        j*                  f      rtG        ||	      S  |	|      S # t,        $ rK}t.        j0                  j2                  r t5        | j6                  j8                  |t;        |            |d }~ww xY w)Nr   r  r  r,  c                 2    t        |       r| S t        |       S rK   )rN   r   )r  s    r-   wrap_constexprz3CodeGenerator.call_Function.<locals>.wrap_constexpr\  s    "Q<r/   )$rL   r  r   insertr  r  r   ri   r  rm   rN   r   rR   
is_builtinr   rG  r  	signature__call__
parametersr#  r   r  r   rL  r  r$   r  ra   rU  r7  rx   mapr   builtinsrw   )
r   rd   re   rf   rN  extra_kwargssigrF   r  r  s
             r-   call_FunctionzCodeGenerator.call_Function4  s   b>+ABCKK2;;'Bb+&4T*((T377B
#(8(E(--JbJbceJfjt%k'6L"/0''4''+cnn,,0MM[)s~~--1\*M$6,6#6c5)"..-C
 ''..00+T2D$#	 
 cHNNHNN;<)#~>>c""-  	M$$88 't{{c!fE1L	Ms   ?0G% %	H9.AH44H9c                 n    t        |t              r|j                  d|       | j                  ||||      S r~   )rL   r   r  r  )r   rd   re   fn_selfrf   rN  s         r-   r  zCodeGenerator.call_Methode  s1    b+&KK7#!!$D#66r/   c                     t         j                  |j                              }t        |t              s& j
                  j                  |      }|	 | |      S t        |dd      }|rt        |dd      rtdt        j                  |j                        z  g}t        |t              r|j                  |       t         j                  j                  |dj                  |            t!         fd|j"                  D              }|j$                  D cg c]  } j                  |       }}t'        t(        j*                  j-                  d |D                    } j/                  ||||      S c c}w )N_must_use_resultF
_is_unusedz#The result of %s is not being used.rJ  c              3   @   K   | ]  }j                  |        y wrK   r   r  s     r-   r   z+CodeGenerator.visit_Call.<locals>.<genexpr>x  s     D74::g&Dr   c              3   F   K   | ]  }t        |t              r|n|g  y wrK   )rL   r[   )r   r  s     r-   r   z+CodeGenerator.visit_Call.<locals>.<genexpr>z  s#     1bXYz!T7J!QRPS2S1bs   !)r   r   r   rL   r   statically_implemented_functionsr  r   r   r7  rU  rk  r$   r  ra   r:   rG  rH  rf   r[   	itertoolschainfrom_iterabler  )	r   rd   re   static_implementationmurerror_messagerN  rh   rf   s	   `        r-   r   zCodeGenerator.visit_Callj  s/   !$**TYY"78"n-$($I$I$M$Mb$Q!$0,T488b,e474u5BS[[QUQZQZE[[\M#s#$$S)"4;;??D#((=:QRRDdmmDD+/995C

355IOO111b]a1bbc!!$D#66 6s   E5c                 ,    t        |j                        S rK   )r   rs   r   s     r-   visit_ConstantzCodeGenerator.visit_Constant~  s    $$r/   c           
         | j                   j                  t        |j                              }|5| j	                  |dj                  |j                  j                              g }|j                  D ]  }| j                  |      }t        |      s't        |      }|du r	|dk(  r|c S |du s<|dk(  sB|c S |j                  j                         rYt        |dd       }||| j                  z  }t        j                  dt         | j"                  |t%        j&                  |             |j)                  |        t+        |      d	k(  r|j)                         t+        |      d
k\  rS|j-                         }|j-                         }	| j/                  ||	|      }
|j)                  |
       t+        |      d
k\  rSt+        |      dk(  sJ |d	   S )Nz9AST boolean operator '{}' is not (currently) implemented.Flogical_andT
logical_orlinenozeLogical operators 'and' and 'or' are deprecated for non-scalar tensors; please use '&' or '|' instead)categoryfilenamer  sourcer   r   r#   )_method_name_for_bool_opr  r+   r  r=  r  rb   rx   r   rP   r   rW   r   r  r5  warn_explicitUserWarningr  r   r7  rk  r   rm  r  )r   rd   r  nontrivial_valuessubnoders   bvr  r  r  rP  s              r-   visit_BoolOpzCodeGenerator.visit_BoolOp  s   3377TWWF##QXXY]Y`Y`YiYijl l {{ 	0G JJw'E$U+%[%Kk]&B L$J[L%@ L ::&&($T8T:F)$//1**!,!%%"{{40 "((/9	0<  !Q& $$U+#$)#'')C#'')C++KcBC$$S)	 #$) $%*** ##r/   r  r  r  c                 v   | j                  |j                        }t        |      r+|j                  dk(  r| j                  j                  |d      S t        |t              r|j                  dvr|j                  }t        ||j                        }t        |      rt        |t              rt        ||      S |S )NT)r#   r   )rs   r+   )r   rs   rP   r   r#  permuterL   r   r   rN   r   r  )r   rd   r  r   s       r-   r   zCodeGenerator.visit_Attribute  s    jj$S!dii3&6==((f55c9%$));L*L))CsDII&C Zk%B!#t,,r/   c                 f    d|j                   _        t        j                  j	                  | |       y r   )rs   r   r   r  r   r   s     r-   
visit_ExprzCodeGenerator.visit_Expr  s"     $

%%dD1r/   c                      y rK   r   r   s     r-   visit_NoneTypezCodeGenerator.visit_NoneType  r   r/   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  )r[   rx   r`   rL   r   ConstantrU  rs   FormattedValue
conversionr   rT   r=  r+   chrr  AssertionErrorr:   )r   rd   rx   rD   rs   conversion_code	evaluateds          r-   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r/   c           	      "   |y t        j                         5  t        j                  dt               t        j                  dt               | j
                  }| j                  j                         }|| _        t        |d      rt        |d      r| j                  j                  | j                  | j                  |j                  z   |j                        }| j                  @| j                  j                  | j                  j!                  | j                  |             n| j                  j                  |       | j                  j                         }	 t"        | I  |      }|r"|| _        | j                  j                  |       |cd d d        S # t&        $ r  t(        $ rU}t*        j,                  j.                  r t'        | j0                  j2                  | j
                  t5        |            d d }~ww xY w# 1 sw Y   y xY w)Nignorer  
col_offset)r5  catch_warningssimplefilterDeprecationWarningPendingDeprecationWarningr/  r   r\  rm   
create_locr  r  r  r-  r%  r&  r[  superr   r$   r  r   rL  r  r  ra   r=   )r   rd   	last_noder'  here_locrF   r  	__class__s          r-   r   zCodeGenerator.visit  s   <$$& 	 !!(,>?!!(,EFI||++-H DMtX&74+F<<224>>4??UYU`U`C`bfbqbqr**6LL(()E)EdF]F]_g)hiLL((2<<//1	ZgmD)  )$$X.=	 	" $  Z$$88 't{{tAwOUYYZ'	 	s1   EHF,%HH-AG==HHHc                 j    | j                  |dj                  t        |      j                              )Nzunsupported AST node type: {})r=  r  r+   rb   r   s     r-   r   zCodeGenerator.generic_visit  s,    &E&L&LTRVZM`M`&abbr/   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: >)r   rf   rH  r  r   r   rL   r   r  r  r=   r%   r  ra   )r   rd   	arg_countpassedr<  r  s         r-   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 )Nrd   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rK   r   r  s     r-   r   z=CodeGenerator.static_executor.<locals>.ret.<locals>.<genexpr>  s     #UGDJJw$7#Ur   )rH  r   rf   r   r   )r   rd   r*   rs   rN  rh   rf   	python_fns   `      r-   rF   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@  rF   s   ` r-   static_executorzCodeGenerator.static_executor  s    	6CHH 	6 
r/   r   )r   r  rK   )rb   r   r   r   r   r   rU  r   r   r[   r  r  r  rL   r   rm   r7  r   r  updater   rR   device_printminimummaximumr=  rA  r2  
contextlibcontextmanagerrX  r^  r   r   r   r`  rd  rf  rz  rl  r   r  r   ListCompr  r   r  r   r  r  r  r  r   r   r   r  r  r  r  r  AddSubMultDivFloorDivModPowLShiftRShiftBitAndBitOrBitXorr  r
   operatorr  r  r*  r   r   rR  rT  r]  EqNotEqLtLtEGtGtErZ  cmpopri  USubUAddNotInvertrg  unaryopro  ro  r  r  r  r  r  r  r  r  r	   r  r  r  r  r  r   r  BoolOpr  AndOrr  boolopr   r  r   r*  r   r   r   r<  rA  experimental.gluonttglstatic_assertstatic_printr8  r  r?  r   __classcell__)r   r1   r6  s   00@r-   r  r  '  s    /3efjQUbc=0+ =0U]^bUc=0AI#=0B tUE3
GWM) 	


A)tCH~  	(--,,-	  !	  ! L	1f ' '
U&c &%
I0E*F &4 &"#-P	2!3<< ! @<S 
6?p&' 2*..
22$	.@ 	)n

L

L

I		8

I>T$s||"4c"9: 7Cr$,N!,F8>3/j*>@$ 	#))Xsvvx(TWTZTZ\dfifmfmow<d4		?C#78 t  	)SXXy#''9cjjR^?tD$5s$:; @-76rL/6A6F2&0U38_ 0OC O
/E; /Eb/#b7
7(%2$ 2$h >AWWmUXU[U[]i<jd4

#3S#89j
2$!Fc#(( t ,
 6##%:""OE$:1?51_S!_S!Q$d68SXXJO3L+L&M {%)s   N6r  c                    d gt        | j                        z  }t        |j                  j	                               }t        |d      \  }	}
t        |j                  j	                               D ]O  \  }\  }}| j                  j                  |      }d }|	|	d   |k(  r|
}t        |d      \  }	}
t        ||      ||<   Q t        g ||j                  |j                        }t        |       \  }}ddlm} t        d |j                        }|D ci c]"  }| j                  |d      |j                  |   $ }}|j                  }  |dddg      ||      }t!        ||| j#                         | j%                  |      | d||||||| j'                         	      }|j)                  | j+                                |j,                  }||_        |j1                         s&| j'                         st3        |       t5        d
      |S c c}w )N)NNr   )
namedtuplec                     t        |       dk(  S rV   )r   )rt   s    r-   r   zast_to_ttir.<locals>.<lambda>?  s    c!fk r/   SpecializationProxyr@   r  T)r   r.  r  r  r  r  r'  r(  r)  r  r   z error encountered during parsing)r   rc   r  r@   r   r  r`   r  indexr   r   r   r   collectionsrl  filterr  r  r=   r   r   r  r  r  verify_with_diagnosticsr8  r  )re   ra   r  r'  r(  r)  r  r   
const_iterkcvcrD   ksrt   rg   cexprr,  r  r  rl  leavesr@   r  proxyr   s                            r-   ast_to_ttirrz  /  s   R\\**Icmm))+,J*l+FB 3 3 56 -
7Bll  $>beqjE*l3FB"1e,	#- B	3==#))DI04Iz&)3==9F?EF!ad#S]]1%55FIFIIJ,{K.HI)U^_Egy9M9M9O_a_f_fgl_m%'49Ycmt*5*U[fhfqfqfsuI OOBHHJFFN))+{{}&M=>>M Gs   'HrK   )Or   r  rF  rq  r  r(   r5  rO  r  dataclassesr   rv   r   typingr   r   r   r   r	   r
   r   r   r   r  r   r   _C.libtritonr   r   r   r   r   r   r  language.corer   r   r   runtime.jitr   r   r   r   r   r   _utilsr    r!   r"   errorsr$   r%   r&   r.   rG   r   rN   rP   rT   rY   r\   ri   ro   rw   r|   rs   r   r  r+   r;  r   r   r   r  r   r   r  r  rz  r   r/   r-   <module>r     s   
     	    !  T T T  ' F F G G C  C H H a a% % %! ! !US UT UNS NT N(S (T (XD*!5 "bhh "Y " #tDz* m3 3*T%COO T%n7 7t $  
ECOO EP(r/   