o
    hG                    @   s,  d dl Z d dlZd dlZd dlZd dlZd dlZd dlZd dlmZ 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 dd	lmZm Z m!Z! dd
l"m#Z#m$Z$ ddl%m&Z& ddl'm(Z(m)Z)m*Z* ddl+m,Z,m-Z-m.Z. dd Z/dd Z0dede1fddZ2dede1fddZ3dede1fddZ4dede1fddZ5dede1fddZ6dd  Z7d!d" Z8d#d$ Z9d%ee  fd&d'Z:d(eej; d)ee! fd*d+Z<e1e=e>dhZ?G d,d- d-Z@G d.d/ d/e jAZBG d0d1 d1ZCed2d3G d4d5 d5ZDG d6d7 d7e jAZEd:d8d9ZFdS );    N)	dataclass)
ModuleType)	AnyCallableDictOptionalTupleTypeUnionIterableList   )knobslanguage)irgluon_ir)	constexpr	str_to_tytensor)_unwrap_if_constexpr
base_value	base_type)get_jit_fn_file_lineget_full_name)JITFunction)find_paths_ifget_iterable_pathset_iterable_path   )CompilationErrorCompileTimeAssertionFailureUnsupportedLanguageConstructc                 C   s*   d}t || std| d|  | | S )Nz^[a-zA-Z_][a-zA-Z0-9_]*$zinvalid z identifier: )rematchr   )nametypepattern r'   Y/var/www/html/scripts/venv/lib/python3.10/site-packages/triton/compiler/code_generator.pycheck_identifier_legality   s   r)   c                    st   d dd |D }d  fddt D }|dd}|dd}|d	dd
d}|  d| d| }|S )N_c                 S      g | ]}|  qS r'   )mangle.0tyr'   r'   r(   
<listcomp>!       zmangle_fn.<locals>.<listcomp>c                    s"   g | ]}| d t  |  qS )c)reprr.   i	constantsr'   r(   r0   "      " ._d_'_sq_[]__)joinsortedreplace)r$   arg_tysr7   mangled_arg_namesmangled_constantsretr'   r6   r(   	mangle_fn   s   rG   oreturnc                 C   
   t | tS N)
isinstancer   rH   r'   r'   r(   _is_triton_value+      
rN   c                 C   rJ   rK   )rL   r   rM   r'   r'   r(   _is_triton_tensor/   rO   rP   c                 C   s   | d u pt | ttjjtfS rK   )rL   r   r   coredtyper   rM   r'   r'   r(   _is_constexpr3   s   rS   c                 C   s   t | o| j o| jjdkS Nr   )rP   r%   is_blocknumelrM   r'   r'   r(   _is_non_scalar_tensor7      rW   c                 C   s   t | ttfS rK   )rL   listtuplerM   r'   r'   r(   _is_list_like;      r[   c              
   C   sX   |j r(t|D ]"\}}t|s't|r't|j| d|j d|j|  d| qd S d S )Nz	Function z= is marked noinline, but was called with non-scalar argument :)noinline	enumeraterS   rW   r!   src__name__	arg_names)nodefnargsidxargr'   r'   r(   _check_fn_args?   s   rh   c                 C   s   t | tot| tot| dS )N_fields)rL   r%   
issubclassrZ   hasattr)valr'   r'   r(   _is_namedtupleI   rX   rm   c                    s~   t t| r
| j}nt| tjr| jj}n	J dt|   fdd| D }dd |D }dd |D }t|t||S )NFzUnsupported type c                       g | ]} |qS r'   r'   r.   vrd   r'   r(   r0   U   r1   z*_apply_to_tuple_values.<locals>.<listcomp>c                 S   s    g | ]}|d u rt |n|qS rK   )r   ro   r'   r'   r(   r0   V   s     c                 S      g | ]}|j qS r'   r%   ro   r'   r'   r(   r0   W       )rm   r%   ri   rL   r   rZ   fields
tuple_type)valuerd   ru   valstypesr'   rq   r(   _apply_to_tuple_valuesM   s   
rz   valuesc                 C   s   g }| D ]}| | q|S rK   )_flatten_ir)r{   handlesrp   r'   r'   r(   flatten_values_to_ir[   s   r~   r}   ry   c                 c   s:    d}|D ]}| | |\}}|V  q|t| ksJ d S )Nr   )_unflatten_irlen)r}   ry   cursorr/   rw   r'   r'   r(   unflatten_ir_valuesb   s   r   c                   @   s$   e Zd Zdd Zdd Zdd ZdS )enter_sub_regionc                 C   
   || _ d S rK   )	generator)selfr   r'   r'   r(   __init__o   rO   zenter_sub_region.__init__c                 C   sL   | j j | _| j j | _i | j _| j j | _| j j	 | _
| j| jfS rK   )r   lscopecopyliveins
local_defs	prev_defsbuilderget_insertion_blockinsert_blockget_insertion_pointinsert_pointr   r'   r'   r(   	__enter__r   s   zenter_sub_region.__enter__c                 O   s(   | j j| j | j| j _| j| j _d S rK   )r   r   restore_insertion_pointr   r   r   r   r   )r   re   kwargsr'   r'   r(   __exit__{   s   
zenter_sub_region.__exit__N)ra   
__module____qualname__r   r   r   r'   r'   r'   r(   r   m   s    	r   c                   @   s  e Zd Zdd ZdefddZdefddZdefdd	Zd
ej	defddZ
d
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZd
ejdefddZdS ) ContainsReturnCheckerc                 C   r   rK   )gscope)r   r   r'   r'   r(   r      rO   zContainsReturnChecker.__init__rI   c                    s   t  fdd|D S )Nc                 3       | ]}  |V  qd S rK   visit)r.   sr   r'   r(   	<genexpr>       z5ContainsReturnChecker._visit_stmts.<locals>.<genexpr>)any)r   bodyr'   r   r(   _visit_stmts   s   z"ContainsReturnChecker._visit_stmtsc                 C      dS NFr'   )r   rd   r'   r'   r(   _visit_function   s   z%ContainsReturnChecker._visit_functionc                 C   sf   d}t |D ])\}}t|tr#|D ]}t|t jr!|p | |}qqt|t jr0|p/| |}q|S r   )astiter_fieldsrL   rY   ASTr   )r   rc   rF   r*   rw   itemr'   r'   r(   generic_visit   s   
z#ContainsReturnChecker.generic_visitrc   c                 C   sP   t |jtjr"|jj| jv r | j|jj }t||j}| |S dS | 	|jS r   )
rL   rw   r   Nameidr   getattrattrr   r   )r   rc   rw   rd   r'   r'   r(   visit_Attribute   s   
z%ContainsReturnChecker.visit_Attributec                 C   s:   t |jtju r
dS |j| jv r| j|j }| |S dS r   )r%   ctxr   Storer   r   r   )r   rc   rd   r'   r'   r(   
visit_Name   s   
z ContainsReturnChecker.visit_Namec                 C   r   NTr'   r   rc   r'   r'   r(   visit_Return      z"ContainsReturnChecker.visit_Returnc                 C   r   r   r'   r   r'   r'   r(   visit_Assign      z"ContainsReturnChecker.visit_Assignc                 C   r   r   r'   r   r'   r'   r(   visit_AugAssign   r   z%ContainsReturnChecker.visit_AugAssignc                 C      |  |jS rK   r   r   r   r'   r'   r(   visit_Module      z"ContainsReturnChecker.visit_Modulec                 C   r   rK   r   r   r'   r'   r(   visit_FunctionDef   r   z'ContainsReturnChecker.visit_FunctionDefc                 C   s&   |  |j}|jr|p|  |j}|S rK   )r   r   orelse)r   rc   rF   r'   r'   r(   visit_If   s   zContainsReturnChecker.visit_Ifc                 C   s   |  |jp|  |jS rK   )r   r   r   r   r'   r'   r(   visit_IfExp      z!ContainsReturnChecker.visit_IfExpc                 C   r   rK   )r   funcr   r'   r'   r(   
visit_Call   r   z ContainsReturnChecker.visit_CallN)ra   r   r   r   boolr   r   r   r   	Attributer   r   r   Returnr   Assignr   	AugAssignr   Moduler   FunctionDefr   Ifr   IfExpr   Callr   r'   r'   r'   r(   r      s    r   c                   @   sh   e Zd Zdd Zdejdee deej fddZ	dejdeej fdd	Z
dejfd
dZdd ZdS )ASTFunctionc                 C   s   || _ || _|| _|| _d S rK   )	ret_types	arg_typesr7   attrs)r   r   r   r7   r   r'   r'   r(   r      s   
zASTFunction.__init__r   ry   rI   c                 C   s(   g }|D ]}|d u rq| || q|S rK   )_flatten_ir_types)r   r   ry   ir_typesr/   r'   r'   r(   flatten_ir_types   s   zASTFunction.flatten_ir_typesc                 C   s   |  || jS rK   )r   r   )r   r   r'   r'   r(   return_types_ir   r\   zASTFunction.return_types_irc                    sP    fdd}t t j|} fdd|D } ||} |}|||S )Nc                       |  j vo|d uS rK   r6   pathr*   r   r'   r(   <lambda>   rt   z'ASTFunction.serialize.<locals>.<lambda>c                    s   g | ]}t  j|qS r'   )r   r   r.   r   r   r'   r(   r0          z)ASTFunction.serialize.<locals>.<listcomp>)rY   r   r   r   r   get_function_ty)r   r   is_val	val_pathsr   arg_types_irret_types_irr'   r   r(   	serialize   s   
zASTFunction.serializec                    s   fddj }fdd}ttj |}d} fddt  D }|D ]+}tj |}j|g }	|	D ]\}
} ||
| q;|	||\}}t
||| q*j}| D ]\}}t
||t| q]|S )Nc                    s6   t | tttjfrt fdd| D | S td S )Nc                    rn   r'   r'   r.   xmake_templater'   r(   r0      r1   zBASTFunction.deserialize.<locals>.make_template.<locals>.<listcomp>)rL   rY   rZ   r   rv   r   )r/   r   r'   r(   r      s   
z.ASTFunction.deserialize.<locals>.make_templatec                    r   rK   r6   r   r   r'   r(   r      rt   z)ASTFunction.deserialize.<locals>.<lambda>r   c                       g | ]}  |qS r'   re   r4   rq   r'   r(   r0         z+ASTFunction.deserialize.<locals>.<listcomp>)r   rY   r   rangeget_num_argsr   r   getset_arg_attrr   r   r7   itemsr   r   )r   rd   rx   r   r   r   r}   r   r/   
attr_specs	attr_nameattr_valrl   r7   r'   )rd   r   r   r(   deserialize   s"   
zASTFunction.deserializeN)ra   r   r   r   r   r   r   r   r%   r   r   r   r   r'   r'   r'   r(   r      s    "
r   T)frozenc                   @   s   e Zd ZU eed< eed< dS )BoundJITMethod__self____func__N)ra   r   r   r   __annotations__r   r'   r'   r'   r(   r     s   
 r   c                       s  e Zd ZU 		ddedee dee fddZd	d
 ee	e
eeeeefD Zeeef ed< edejjfdejfdejff dd Zdd Zdd Zdedeeef ddfddZdd Z dd Z!dd Z"d d! Z#d"d# Z$d$d% Z%defd&d'Z&d(d) Z'd*d+ Z(d,d- Z)d.d/ Z*d0d1 Z+d2d3 Z,d4d5 Z-d6d7 Z.d8d9 Z/d:d; Z0d<d= Z1d>d? Z2d@dA Z3e4j5dBe4j6dCe4j7dDe4j8dEe4j9dFe4j:dGe4j;dHe4j<dIe4j=dJe4j>dKe4j?dLe4j@dMiZAeeBe4jC ef edN< dOdP ZDdQdR ZEdSdT ZFdUdV ZGdWdX ZHdYdZ ZId[d\ ZJe4jKd]e4jLd^e4jMd_e4jNd`e4jOdae4jPdbiZQeeBe4jR ef edc< ddde ZSe4jTdfe4jUdge4jVdhe4jWdiiZXeeBe4jY ef edj< dkdl ZZdmdn Z[dodp Z\dqdr Z]dsdt Z^dudv Z_dwdx Z`dydz Zad{d| Zbd}d~ Zcdd Zddeeeef fddZfdefddZgdefddZhdd Zidd Zjde4jkfddZle4jmde4jndiZoeeBe4jp ef ed< dd Zqdd Zrdd Zsdd Zt fddZudd Zvde4jwddfddZxdd ZyddlzmZ{ ejj|exejj}eye~e{j|exe{j}eye~eeyeeeyeiZeeee4jwgef f ed<   ZS )CodeGeneratorNFr   jit_fnfunction_types	file_namec                 C   s  || _ | rddlm} t|| _|| j| _nddlm	} t
|| _|| j| _|| _|d | _| j||d || j_|| j_|d u rJi n|| j_|	d u rW| j n|	| _|d u r`i n|| _|| _i | _| D ]/\}}t|tr||j|| j|< qmt|dd}||v rt|| |j| j|< qm|| j|< qmi | _|| _|
r||dd d  }t|d}|| _ |
| _!d | _"|| _#g | _$d | _%i | _&| ' | _(d | _)d	| _*d S )
Nr   )GluonSemantic)TritonSemanticr   r    r9   functionF)+contextis_gluon,triton.experimental.gluon.language._semanticr   r   GluonOpBuilderr   semantictriton.language.semanticr   r   r   
begin_lineset_locoptionscodegen_fns
module_mapcreate_modulemodulefunction_ret_types	prototyper   r   rL   r   r   ra   r   r   r   rfindr)   function_name	is_kernelcur_noder^   	scf_stackret_typer   _define_name_lookupdereference_namerd   visiting_arg_default_value)r   r  r  r   r  r   r	  r
  r  r  r  r   r^   r   r  r   r   krp   module_namer'   r'   r(   r     sR   




zCodeGenerator.__init__c                 C   s   i | ]}|j |qS r'   ra   r.   r*   r'   r'   r(   
<dictcomp>V  s    zCodeGenerator.<dictcomp>builtin_namespaceprintminmaxc                 C   s   t | jj||S rK   )r!   r   r`   )r   rc   messager'   r'   r(   _unsupported`  s   zCodeGenerator._unsupportedc                 C   s0   t  }| j||}||u rdS t|rdS dS )NFT)objectr   r   rS   )r   r$   absent_markerrl   r'   r'   r(   _is_constexpr_globalc  s   z"CodeGenerator._is_constexpr_globalc                    sJ   dt ffdddt ffddt  dt dtf fdd}|S )	Nr$   c                    s    j | |S rK   )r   r   )r$   absentr   r'   r(   local_lookupp  s   z7CodeGenerator._define_name_lookup.<locals>.local_lookupc                    s    j | |}t||u |  jv t|tu t|tt|ddt|ddt|dd	dt|dd	dt|t
jt| |  jtjjgrH|S ttd|  d	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   r  r%   r   rL   r   r   
startswithr   rR   rm   r&  r  r   compilationallow_non_constexpr_globals	NameErrortextwrapdedentrB   )r$   r'  rl   r   r'   r(   global_lookupt  s.   



z8CodeGenerator._define_name_lookup.<locals>.global_lookuprI   c                    s@    }j jfD ]}|| |}||ur|  S q	t|  d)Nz is not defined)r  r   r0  )r$   r'  lookup_functionrw   r%  r3  r(  r   r'   r(   name_lookup  s   
z6CodeGenerator._define_name_lookup.<locals>.name_lookup)strr$  r   )r   r6  r'   r5  r(   r  n  s
   !z!CodeGenerator._define_name_lookupr$   rw   rI   c                 C   s   || j |< || j|< dS )z This function:
            called by visit_Assign() & visit_FunctionDef() to store left value (lvalue)
        1. record local defined name (FIXME: should consider control flow)
        2. store tensor in self.lvalue
        N)r   r   )r   r$   rw   r'   r'   r(   	set_value  s   
zCodeGenerator.set_valuec                 C   s   | j  }| j  }||fS rK   )r   get_locr   )r   locipr'   r'   r(   _get_insertion_point_and_loc  s   

z*CodeGenerator._get_insertion_point_and_locc                 C   s   | j | | j | d S rK   )r   r   r  )r   r;  r:  r'   r'   r(   _set_insertion_point_and_loc  s   z*CodeGenerator._set_insertion_point_and_locc                 C   s8   t |s|g}|D ]}| | t|tjr d S q	d S rK   )r[   r   rL   r   r   )r   stmtsstmtr'   r'   r(   visit_compound_statement  s   
z&CodeGenerator.visit_compound_statementc                 C      t j| | d S rK   r   NodeVisitorr   r   r'   r'   r(   r        zCodeGenerator.visit_Modulec                    s6     |j}|d u sJ t fdd|jD }|S )Nc                    r   r'   r   )r.   eltr   r'   r(   r0     r   z,CodeGenerator.visit_List.<locals>.<listcomp>)r   r   r   rZ   elts)r   rc   r   rF  r'   r   r(   
visit_List  s   zCodeGenerator.visit_Listc                    s    |j}g } fdd  |}|d u rtj}nt|tjjs$J || |j}j	
| jd u r;|_nj|krKtdj d| j	 }j	| d S )Nc                    s8   t | tjrt|  S t | tjttfrj| S | S rK   )	rL   r   rZ   rz   r   intfloatr  	to_tensor)rw   decayr   r'   r(   rL    s
   
z)CodeGenerator.visit_Return.<locals>.decayzInconsistent return types:  and )r   rw   r   voidrL   rQ   r   r|   r%   r   rF   r  	TypeErrorcreate_blockset_insertion_point_to_end)r   rc   	ret_valuer}   ret_typost_ret_blockr'   rK  r(   r     s    



zCodeGenerator.visit_Returnc                 C   s$   |  |j}t|tjjsJ |jS rK   )r   rw   rL   r   rQ   rZ   r{   r   rc   re   r'   r'   r(   visit_Starred  s   zCodeGenerator.visit_Starredc              	      s,    |j\}} jr |dt|jjd d d D ]G\}}|jj| d  }|j}|j}tj	|t
 d}	|d u rEtj|	g|d}
ntj|	||d}
z jrSJ d _  |
 W d _qd _w  jrid	nd
} j j} j j j|| j _ j j  j } j j}t||D ]
\}} || q j } j|  |j  j  rJ  j d u sȈ j t!j"krt!j" _  j#g  n1t$ j t!j%r j j& j_'n j g j_' j( j j  j# fdd j) jD   j*  |r j+| d S d S )Nz,nested function definition is not supported.r   )r   r   targetsrw   )targetrw   
annotationTFpublicprivatec                    s   g | ]} j |qS r'   )r   create_poisonr-   r   r'   r(   r0      r   z3CodeGenerator.visit_FunctionDef.<locals>.<listcomp>),r   re   rd   r#  r_   defaultsr[  rg   r   r   r   r   	AnnAssignr  r  r  r   r   get_or_insert_functionr  r  r^   	push_backadd_entry_blockr   zipr8  r   set_insertion_point_to_startr@  r   has_terminatorr  r   rN  rF   rL   rv   ry   r   
reset_typer   finalizerQ  )r   rc   rb   kwarg_namesr5   default_valuearg_noder[  r$   	st_target	init_node
visibilityfn_tyentry
arg_valuesarg_name	arg_value	insert_ptr'   r   r(   r     sP   


$
zCodeGenerator.visit_FunctionDefc                 C   s4   g }|j D ]
}|| |g7 }q| |j}||fS rK   )re   r   kwarg)r   rc   rb   rg   ri  r'   r'   r(   visit_arguments&  s
   
zCodeGenerator.visit_argumentsc                 C   s   t j| | |jS rK   )r   rC  r   rg   r   r'   r'   r(   	visit_arg-  s   zCodeGenerator.visit_argc                 C   sj   |  |j}|  |j}|  |j}|tkr0|| jv r"t| dt|}|| j|< | j| S | |S )Nz4 is already defined. constexpr cannot be reassigned.)r   r[  rZ  rw   r   r   
ValueErrorr   )r   rc   r[  rZ  rw   r'   r'   r(   visit_AnnAssign1  s   



zCodeGenerator.visit_AnnAssignc                 C   s   t |jtjs	J t |tjr| ||S t |tjr0t|jD ]\}}| 	||j
|  q d S t |tjrE| |j}t||j| d S t |tjsMJ | | || d S rK   )rL   r   r   r   	Subscriptvisit_Subscript_Storer   r_   rF  assignTargetr{   r   r   rw   setattrr   r   r8  )r   rZ  rw   r5   baser'   r'   r(   r|  A  s   zCodeGenerator.assignTargetc                    s\    fdd   |j}t|tjr|jgn|j}t|dks$J |d | d S )Nc                    sT   t | tjrt|  S tjtjf}t| } | d ur(t| s(t | |s(j| } | S rK   )	rL   r   rZ   rz   rR   r   rN   r  rJ  )rw   native_nontensor_types_sanitize_valuer   r'   r(   r  R  s   
z3CodeGenerator.visit_Assign.<locals>._sanitize_valuer   r   )	r   rw   rL   r   r`  rZ  rY  r   r|  )r   rc   r{   rY  r'   r  r(   r   P  s
   zCodeGenerator.visit_Assignc                 C   sN   t |j}t |_t||j|j}tj	|jg|d}| 
| | 
|S )NrX  )r   deepcopyrZ  r   Loadr   BinOpoprw   r   r   )r   rc   lhsrhsassignr'   r'   r(   r   b  s   


zCodeGenerator.visit_AugAssignc                 C   s"   t |jtju r|jS | |jS rK   )r%   r   r   r   r   r  r   r'   r'   r(   r   j  s   zCodeGenerator.visit_Namec                 C   rA  rK   rB  r   r'   r'   r(   visit_Storeo  rD  zCodeGenerator.visit_Storec                 C   rA  rK   rB  r   r'   r'   r(   
visit_Loadr  rD  zCodeGenerator.visit_Loadc                    s    fdd|j D }t|S )Nc                    r   r'   r   r   r   r'   r(   r0   v  r   z-CodeGenerator.visit_Tuple.<locals>.<listcomp>)rF  r   rZ   rU  r'   r   r(   visit_Tupleu  s   
zCodeGenerator.visit_Tuplec                 C   sv   t |rt|||| jdS t |r#tdd|}t|||| jdS t|ttjfs4t|tr4t|}t|||S )N	_semanticz__(.*)__z__r\1__)	rP   r   r  r"   subrL   r   r   rZ   )r   method_namer  r  reverse_method_namer'   r'   r(   _apply_binary_methody  s   z"CodeGenerator._apply_binary_methodc                 C   sV   |  |j}|  |j}| jt|j}|d u r$| |d|jj	| 
|||S )Nz8AST binary operator '{}' is not (currently) implemented.)r   leftright_method_name_for_bin_opr   r%   r  r#  formatra   r  )r   rc   r  r  r  r'   r'   r(   visit_BinOp  s   zCodeGenerator.visit_BinOp__add____sub____mul____truediv____floordiv____mod____pow__
__lshift__
__rshift____and____or____xor__r  c                 C   s  | j | | |j | j  }| j }i }|jr9| j | | | _i | _| |j | j }| j  }g }|D ]h}|df|dffD ]4\}	}
||	v r{t	|	| t	|| k}|rg|	| j	|| j	ks{J d| d||  d|
 d|	|  qG||v s||v r|
| ||v r||vr|| ||< ||v r||vr|| ||< q=t| | @ D ]6}||v rq|| }|j	}|| }|j	}t	|t	|k}|r||ksJ d| d| d	| d
|
| q|||||fS )Nthenelsezinitial value for `z` is of type z
, but the z block redefines it as zMismatched type for z between then block (z) and else block ())r   re  r@  r   r   r   r   r   r   r%   appendrA   keys)r   rc   r   
then_block
else_block	then_defs	else_defsnamesr$   defs
block_name
type_equalthen_valthen_tyelse_valelse_tyr'   r'   r(   visit_then_else_blocks  s^   





z$CodeGenerator.visit_then_else_blocksc                    s  t | }|\}}| j }| j }| j| | j|j|| | ||||\ }}}| j | j| | rCJ | tfdd|D }	| j	|	 | j| | rdJ | t fdd|D }
| j	|
 t
|	t
|
ksJ t|	|
D ]\}}| }|| ksJ | qW d    n1 sw   Y  | j fddtt
|	D }fdd|D }t||}t||D ]
\}}| || qd S )Nc                 3       | ]} | V  qd S rK   r'   r.   r$   r  r'   r(   r         z3CodeGenerator.visit_if_top_level.<locals>.<genexpr>c                 3   r  rK   r'   r  r  r'   r(   r     r  c                    r   r'   rg   r4   )endif_blockr'   r(   r0     r   z4CodeGenerator.visit_if_top_level.<locals>.<listcomp>c                       g | ]} | j qS r'   rs   r  r  r'   r(   r0     r   )r   r   rP  rQ  create_cond_branchhandler  rf  r~   create_branchr   rd  get_typeadd_argumentre  r   r   r8  )r   condrc   srr   ip_blockr  r  r  then_handleselse_handlesthen_helse_hr/   res_handlesry   
new_valuesr$   	new_valuer'   )r  r  r  r(   visit_if_top_level  s>   




z CodeGenerator.visit_if_top_levelc                    s  t | }|\}}|  \}}| j }|jr| j nd }	| ||||	\ }}	}
tfdd|
D }| || | jdd |D |j	d|
  | j  t|
dkre| j| |jsm }	n|	
  | j  t|
dkrt fdd|
D }| j| W d    n1 sw   Y  fddtt|D }fd	d|
D }t||}t|
|D ]
\}}| || qd S )
Nc                 3   r  rK   r'   r  r  r'   r(   r      r  z-CodeGenerator.visit_if_scf.<locals>.<genexpr>c                 S   r+   r'   r  r.   hr'   r'   r(   r0     r1   z.CodeGenerator.visit_if_scf.<locals>.<listcomp>Tr   c                 3   r  rK   r'   r  r  r'   r(   r     r  c                    r   r'   
get_resultr4   )if_opr'   r(   r0     r   c                    r  r'   rs   r  r  r'   r(   r0     r   )r   r<  r   rP  r   r  r~   r=  create_if_opr  merge_block_beforeget_then_blockrQ  r   create_yield_opget_else_blockr   r   rd  r8  )r   r  rc   r  r   r*   r;  last_locr  r  r  r  r  r  ry   r  r$   r  r'   )r  r  r  r(   visit_if_scf  s:   



zCodeGenerator.visit_if_scfc              	   C   s  |  |j}t|rZt|r| |d|j r.tdt	
|j  tjj|| j| d}|jtj| jd}t| j |}|rR| jrJ| |d| || d S | || d S t|}t|tvry| |dddd	 tD t|j|r~|jn|j}| | d S )
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  zCannot have `return` statements inside `while` or `for` statements in triton (note that this also applies to `return` statements that are inside functions transitively called from within `while`/`for` statements)O`if` conditionals can only accept values of type {{{}}}, not objects of type {}, c                 s       | ]}|j V  qd S rK   r  r  r'   r'   r(   r   2      z)CodeGenerator.visit_If.<locals>.<genexpr>)r   testrP   rW   r#  r%   rU   warningswarnr   unparser   rQ   _unsplatr  toint1r   r   r  r  r  r   _condition_typesr  r@   ra   r   r   r@  )r   rc   r  contains_returnactive_blockr'   r'   r(   r     s<   

zCodeGenerator.visit_Ifc              	   C   s  |  |j}t|r|jtj| jd}t|  |  \}}| j	
 }| j	| | j|  |j}| j	 }| j	
 }| j	| | j|  |j}| j	 }| || |j|jksjJ d|j d|j |j}	|	tjkry|	| j	gng }
| j	|
|jd}||  |
r| j	|  | j	|jg | j	|  ||  |
r| j	|  | j	|jg |
rtj|d|	nd W  d    S 1 sw   Y  d S t|}t|tvr|  |d!d"dd	 tD t|j#|r|  |jS |  |jS )
Nr  zATernary expression with dynamic condition has inconsistent types rM  Tr   r  r  c                 s   r  rK   r  r  r'   r'   r(   r   f  r  z,CodeGenerator.visit_IfExp.<locals>.<genexpr>)$r   r  rP   r  r   r  r  r   r<  r   rP  re  rJ  r   r   r   r=  r%   rN  to_irr  r  r  r  rQ  r  r  rQ   r   r  r   r  r#  r  r@   ra   )r   rc   r  r;  r  r  r  r  r  r  ret_type_irr  r'   r'   r(   r   8  sT   




$#zCodeGenerator.visit_IfExpc                 C      d S rK   r'   r   r'   r'   r(   
visit_Passm  r   zCodeGenerator.visit_Passc                 C   s   t |jdkrt |jdks| |d| |j}| |jd }t|}t|}t|jd tj	u r:t
||u S t|jd tju rJt
||uS | jt|jd }|d u rf| |d|jd j| |||S )Nr   z1simultaneous multiple comparison is not supportedr   z<AST comparison operator '{}' is not (currently) implemented.)r   comparatorsopsr#  r   r  r   r%   r   Isr   IsNot_method_name_for_comp_opr   r  ra   r  )r   rc   r  r  	lhs_value	rhs_valuer  r'   r'   r(   visit_Comparep  s    zCodeGenerator.visit_Compare__eq____ne____lt____le____gt____ge__r  c                 C   s   |  |j}| jt|j}|d u r| |d|jj dt|r,t	||| j
dS zt	|| W S  tyT   |dkrEt|  Y S | |d| dt|j 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#  ra   rP   r   r  AttributeErrorr   )r   rc   r  rd   r'   r'   r(   visit_UnaryOp  s   zCodeGenerator.visit_UnaryOp__neg____pos__r  
__invert__r  c                 C   s   t |sJ d| dt |sJ d| dt|t|u s(J d| dt|rB|j|jksDJ d| d|j d|j d	d S d S )
Nzcannot reassign constxpr z in the loopzcannot reasign constexpr zLoop carried variable z changed typezLoop-carried variable z has initial type z but is re-assigned to z: in loop! Please make sure that the type stays consistent.)rN   r%   rP   )r   r$   loop_vallive_valr'   r'   r(   _verify_loop_carried_variable  s    z+CodeGenerator._verify_loop_carried_variablec                 C   r   rK   )r   context_exprr   r'   r'   r(   visit_withitem  r   zCodeGenerator.visit_withitemc                    s   t |jdks	J |jd j} |j}|tjkrE fdd|jD }||d ji  	|j
 W d    d S 1 s>w   Y  d S  	|j
 d S )Nr   r   c                    r   r'   r   r.   rg   r   r'   r(   r0     r   z,CodeGenerator.visit_With.<locals>.<listcomp>_builder)r   r   r  r   r   r   
async_taskre   r   r@  r   )r   rc   r  withitemClassre   r'   r   r(   
visit_With  s   
"zCodeGenerator.visit_Withc                    s  t | &}|\}}|  \}}| j }| j| | j| | |j | j	  | j
}|  g }	g }
|D ]}||v rX|| }|| }| ||| |	| |
| q9t|
}dd |D }dd |
D }| || | j||| j || j fddtt|D }t||}t|	|D ]\}}|| j|< || j
|< q| |j}| j | j|j| | j | | j   fddtt|D }t||}t|	|D ]\}}|| j|< || j
|< q| j| | |j | j	  | j
}g }|D ]}||v r|| | q| j| W d    n	1 s/w   Y  fddtt|D }t||}t|	|D ]\}}|| j|< || j
|< qK|jD ]}J dd S )	Nc                 S   r+   r'   r  r  r'   r'   r(   r0     r1   z-CodeGenerator.visit_While.<locals>.<listcomp>c                 S   rr   r'   rs   )r.   ar'   r'   r(   r0     rt   c                    r   r'   r  r4   )before_blockr'   r(   r0     r   c                    r   r'   r  r4   )after_blockr'   r(   r0     r   c                    r   r'   r  r4   )while_opr'   r(   r0     r   FzNot implemented)#r   r<  r   rP  re  r  r  r@  r   popr   eraser  r~   r=  create_while_opcreate_block_with_parent
get_beforer   r   r   rd  r   r   r  rQ  create_condition_opr  	get_afterr|   r  r   r   rC  r   )r   rc   r  r   r   r;  r  dummy	loop_defsr  	init_argsr$   r  r  init_handlesinit_tysinit_fe_tys
block_argscondition_argsrl   r  body_handles	body_argsyieldsresult_handlesresult_valsnew_defr?  r'   )r  r  r  r(   visit_While  sz   









C


zCodeGenerator.visit_Whilec                 C   sJ   t |jtjs	J | |j}| |j}t|r!|j|| j	dS || S )Nr  )
rL   r   r   r  r   rw   slicerP   __getitem__r  )r   rc   r  slicesr'   r'   r(   visit_Subscript_Load  s   z"CodeGenerator.visit_Subscript_Loadc                 C   sJ   t |jtjs	J | |j}| |j}t |tjsJ |	|| d S rK   )
rL   r   r   r   r   rw   r'  r   rZ   __setitem__)r   rc   rw   r  r)  r'   r'   r(   r{    s
   z#CodeGenerator.visit_Subscript_Storec                 C   s
   |  |S rK   )r*  r   r'   r'   r(   visit_Subscript  rO   zCodeGenerator.visit_Subscriptc                    s    fdd|j D S )Nc                    r   r'   r   )r.   dimr   r'   r(   r0     r   z0CodeGenerator.visit_ExtSlice.<locals>.<listcomp>)dimsr   r'   r   r(   visit_ExtSlice  s   zCodeGenerator.visit_ExtSlicec           ,         sN   |jj}fdd|jjD }tfdd|jjD }|tjkrZ||i |}t|j	j
|jj
|jj
}|D ]}t|j|jj< |j |jD ]	}tj| qMq9d S d }	d }
d}d}d}|tju r||i |}|j	}|j}|j}|j}	|j}
|j}|j}|j}n@|tu rt|dkr|d n td}t|dkr|d n |jjd }t|dkr|d n td}ntd	d}t|r|j
dk rt|j
 }d
}||}}j !|}j !|}j !|}|j"# r|j"# r|j"# st$d|j" d|j" d|j" dj %|j"|j"}j %||j"}|&j'}|j(tj)j"j*j+k}|j,}|j,}|j,}j'-|||}j'-|||}j'-|||}j'.|}/|jjtj)0|| t1}|\}}2 \}}j'3 }j'4| j56| |j j57  |8  g }g }g }j9D ]'}||v rɈj9| } || }!:|| |! |6| |6|! |6|  q;|| t<|}"dd |D }#j'=||||" t>|	d ur ?dj'@|	 t>|
d ur ?dj'@|
 |r ?dj'A  |r ?dj'A  |r+ ?dj'A  j56|  Bdj'4 |C _i _9fddtt|"D }$tD|$|#}%tE||%D ]\}}&/||& q[|j j57  g }j9D ]}||v rj9| }'tF|'trj !|'}'|6|' qwt|dkrt<|}(j'G|( H })|)I dksJ dj'4  J }|rшj'K||}j'L||}j|jj j,M| /|jjtj)0|| W d    n	1 sw   Y   fddtt|"D }*tD|*|#}+tE||+D ]\}}&/||& q|jD ]}J dd S )Nc                    r   r'   r   r  r   r'   r(   r0     r   z+CodeGenerator.visit_For.<locals>.<listcomp>c                 3   r   rK   r   r.   keywordr   r'   r(   r     r   z*CodeGenerator.visit_For.<locals>.<genexpr>Fr   r   r   zAOnly `range` and `static_range` iterators are currently supportedTz0For loop bounds and step must all be ints, are (r  r  c                 S   rr   r'   rs   ro   r'   r'   r(   r0   u  rt   ztt.num_stagesztt.loop_unroll_factorztt.disallow_acc_multi_bufferz
tt.flattenztt.warp_specializec                    s   g | ]	}  |d  qS )r   r  r4   )for_op_bodyr'   r(   r0         z7We use SCF, so the loop body should only have one blockc                    r   r'   r  r4   )for_opr'   r(   r0     r   z)Don't know what to do with else after for)Nr   iterr   re   dictkeywordsr   static_ranger   startrw   endstepr   r   rZ  r   r@  r   r   r   rC  r   
num_stagesloop_unroll_factordisallow_acc_multi_bufferflattenwarp_specializer   NumRuntimeErrorrS   r  rJ  rR   is_intrO  integer_promote_implr  r   int_signednessrQ   
SIGNEDNESSSIGNEDr  create_int_castr^  r8  r   r   r<  rP  re  r  r  r  r  r   r  r=  r~   create_for_opr   set_attrget_int32_attrget_unit_attrget_bodyr   r   rd  rL   r  
get_parentsizeget_induction_var
create_sub
create_addreplace_all_uses_with),r   rc   IteratorClass	iter_argsiter_kwargsiteratorr8  r5   r?  r<  r=  r>  r?  r@  lbubr;  negative_stepiv_type
iv_ir_typeiv_is_signedivr  r   r   r;  r  blockr  r"  r  r$   r  r  r  r  block_handlesr  rl   localyield_handlesfor_op_regionr#  result_valuesr'   )r4  r2  r   r(   	visit_For  s  


$&&
$"















Q

zCodeGenerator.visit_Forc                 C   s2   |  |j}|  |j}|  |j}t|||S rK   )r   lowerupperr;  r   r'  )r   rc   rf  rg  r;  r'   r'   r(   visit_Slice  s   zCodeGenerator.visit_Slicec                 C   r   rK   )r   rw   r   r'   r'   r(   visit_Index  r   zCodeGenerator.visit_Indexc                 C   s   |j | |jfS rK   )rg   r   rw   r   r'   r'   r(   visit_keyword  rD  zCodeGenerator.visit_keywordc                 C   s:   |  |j}|jd ur|  |jnd}tjj||| jdS )Nr   r  )r   r  msgr   rQ   device_assertr  )r   rc   r  rk  r'   r'   r(   visit_Assert  s   zCodeGenerator.visit_Assertrd   c                    s  t j|jg R i |  fdd|jD  t D ]\}}t|tjtt	t
tfr2tj| |< qt dd } fdd|D }t dd } fdd|D }tt|d	d |D |}	| j|	st|\}
}d
d  D }tg ||t }t| j|| | j||	| j|j|
|| jj| jj| jjd}z	||   W n t!y } zt"j#j$r t%| j&j'| j(d |d }~ww |j)}|| j|	< n| j|	 }| j*|	}t+|}| j,|||tj-krd S fddt./ D }t0t1||gS )Nc                    s   g | ]} | qS r'   r'   r  r   r'   r(   r0     r1   z2CodeGenerator.call_JitFunction.<locals>.<listcomp>c                 S   s   t |S rK   rS   r*   r   r'   r'   r(   r     s    z0CodeGenerator.call_JitFunction.<locals>.<lambda>c                    s   i | ]}|t  |qS r'   r   r   r   r'   r(   r    r   z2CodeGenerator.call_JitFunction.<locals>.<dictcomp>c                 S   s
   t | S rK   rn  ro  r'   r'   r(   r     s   
 c                    s   g | ]}t  |qS r'   rp  r   r   r'   r(   r0     r   c                 S   rr   r'   rs   r  r'   r'   r(   r0     rt   c                 S   s6   g | ]}|d u st |tttjjfrtjjn|jqS rK   )rL   r   rH  r   rQ   rR   r   r%   r  r'   r'   r(   r0     s    )
r  r   r  r   r^   r   r  r	  r
  r  c                    r   r'   r  r4   )call_opr'   r(   r0     r   )2inspectgetcallargsrd   rb   r_   rL   r   rR   rI  rH  r   r   rQ   r   r   rG   r   r  has_functionr   r   r6  r   r  get_capture_scoper  r^   r   r	  r
  r  r   parse	Exceptionr   r.  front_end_debuggingr   r   r`   r  r  get_functionr~   callrN  r   get_num_resultsnextr   )r   rd   re   r   r5   rg   args_cst	args_pathargs_valfn_namer   r  r   r  r   ecallee_ret_typesymbolr}   r'   )re   rq  r(   call_JitFunction  sR   

zCodeGenerator.call_JitFunctionc              
      s  t  |j}t|ts j|}|d ur| |S t|dd}|rHt|ddrHdt	|j g}t|t
r=|| t jj|d|t fdd|jD } fdd	|jD }ttjd
d |D }t|trz|d|j |j}t|trt|||  |||S t|drt|jstj !|rd j"i}t#$|}	d|	j%v r |d< z||i ||}
t|
t&rt&|
}
|
W S  t'y } zt(j)j*r҂ t jj|d |d }~ww | j+, v rt-t |}||i |}
t.t/|
rt0|
dd S |
S )N_must_use_resultF
_is_unusedz#The result of %s is not being used.r,  c                 3   r   rK   r   r0  r   r'   r(   r     r   z+CodeGenerator.visit_Call.<locals>.<genexpr>c                    r   r'   r   r  r   r'   r(   r0     r   z,CodeGenerator.visit_Call.<locals>.<listcomp>c                 s   s$    | ]}t |tr|n|gV  qd S rK   )rL   rY   r   r'   r'   r(   r      s   " r   r   r  r  c                 S   s   | S rK   r'   )r   r'   r'   r(   r      s    z*CodeGenerator.visit_Call.<locals>.<lambda>)1r   r   r   rL   r    statically_implemented_functionsr   r   r   r  r7  r  r   r   r`   r@   r6  r7  re   rY   	itertoolschainfrom_iterableinsertr   r   r   rh   r  rk   rN   r   rQ   
is_builtinr  rr  	signature
parametersrZ   rw  r   r.  rx  r  r{   maprm   r%   rz   )r   rc   rd   static_implementationmurerror_messagekwsre   extra_kwargssigrF   r  r'   r   r(   r     sP   





 





 zCodeGenerator.visit_Callc                 C   s
   t |jS rK   )r   rw   r   r'   r'   r(   visit_Constant"  rO   zCodeGenerator.visit_Constantrc   c              	   C   sH  | j t|j}|d u r| |d|jjg }|jD ]O}| |}t	|sEt
|}|du r8|dkr8|  S |du rD|dkrD|  S q|j rgt|dd }|d urY|| j7 }tjdt| j|t|d || qt|d	krx|| t|d
kr| }| }	| ||	|}
||
 t|d
ks~t|dksJ |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  ra   r{   r   rP   r   rU   r   r  r  warn_explicitUserWarningr   r   r  r  r   r  r  )r   rc   r  nontrivial_valuessubnoderw   bvr  r  r  resr'   r'   r(   visit_BoolOp%  sJ   





zCodeGenerator.visit_BoolOpr  r  r  c                 C   sr   |  |j}t|r|jdkr| j|dS t|tr#|jdkr#|j}t||j}t	|r7t|t
r7t||S |S )NT)r   r   rw   )r   rw   rP   r   r  permuterL   r   r   rN   r   r   )r   rc   r  r   r'   r'   r(   r   [  s   
zCodeGenerator.visit_Attributec                 C   s   d|j _tj| | d S r   )rw   r  r   rC  r   r   r'   r'   r(   
visit_Exprg  s   zCodeGenerator.visit_Exprc                 C   r  rK   r'   r   r'   r'   r(   visit_NoneTypek  r   zCodeGenerator.visit_NoneTypec                 C   s   t |j}t|D ]N\}}t|tjrt|j||< q	t|tjrO|j	}| 
|j}t|s:| |dtt| |dk r@dndt| d |j||< q	tdt|d|S )Nz^Cannot evaluate f-string containing non-constexpr conversion values, found conversion of type r   z{}z{!}z:encountered unexpected node of type {} in a JoinedStr noder   )rY   r{   r_   rL   r   Constantr7  rw   FormattedValue
conversionr   rS   r#  r%   chrr  AssertionErrorr@   )r   rc   r{   r5   rw   conversion_code	evaluatedr'   r'   r(   visit_JoinedStrn  s"   

*
zCodeGenerator.visit_JoinedStrc                    s  |d u rd S t  v t dt t dt | j}| j }|| _t|dr?t|dr?| j	| j
| j|j |j | j }zt |}W n$ tyO     tyk } ztjjr[ t| jj| jt|d d }~ww |rw|| _| j	| |W  d    S 1 sw   Y  d S )Nignorer  
col_offset)r  catch_warningssimplefilterDeprecationWarningPendingDeprecationWarningr  r   r9  rk   r  r   r  r  r  superr   r   rw  r   r.  rx  r   r`   r3   )r   rc   	last_noder  rF   r  	__class__r'   r(   r     s4   


$zCodeGenerator.visitc                 C   s   |  |dt|j)Nzunsupported AST node type: {})r#  r  r%   ra   r   r'   r'   r(   r     r   zCodeGenerator.generic_visitc              
   C   s   t |j}d|  k rdkrn tdt |jrtdt| |jd }t|ts0td|sh|dkr9d}n%z
| |jd }W n t	y] } zdt
| d }W Y d }~nd }~ww t| jj|t|d S )	Nr   r   z=`static_assert` requires one or two positional arguments onlyzqAssertion condition could not be determined at compile-time. Make sure that it depends only on `constexpr` valuesr   r   z'<failed to evaluate assertion message: >)r   re   r7  rO  r   r   rL   r   NotImplementedErrorrw  r3   r    r   r`   )r   rc   	arg_countpassedr"  r  r'   r'   r(   execute_static_assert  s*   


z#CodeGenerator.execute_static_assertc                    s   dt jf fdd}|S )Nrc   c                    sD   dd  fdd|j D D } fdd|jD }t|i |S )Nc                 S   s   i | ]	\}}|t |qS r'   )r   )r.   r$   rw   r'   r'   r(   r    s    z>CodeGenerator.static_executor.<locals>.ret.<locals>.<dictcomp>c                 3   r   rK   r   r0  r   r'   r(   r     r   z=CodeGenerator.static_executor.<locals>.ret.<locals>.<genexpr>c                    s   g | ]	}t  |qS r'   )r   r   r  r   r'   r(   r0     r3  z>CodeGenerator.static_executor.<locals>.ret.<locals>.<listcomp>)r7  re   r   )r   rc   r  re   	python_fnr   r(   rF     s
   z*CodeGenerator.static_executor.<locals>.ret)r   r   )r  rF   r'   r  r(   static_executor  s   zCodeGenerator.static_executorr   )r   r  )NFNFNr   )ra   r   r   r   r   r   r7  r   r   rY   r   rI  rH  rL   r   rk   r  r   r   updater   rQ   device_printminimummaximumr#  r&  r  r
   r   r   r8  r<  r=  r@  r   rG  r   rV  r   rv  rw  ry  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   r  r  EqNotEqLtLtEGtGtEr  cmpopr  USubUAddNotInvertr  unaryopr  r  r  r&  r*  r{  r,  r/  re  rh  ri  r   rj  rm  r  r   r  BoolOpr  AndOrr  boolopr   r  r  r  r   r   r   r  r  experimental.gluonttglstatic_assertstatic_printr  r  r$  r   __classcell__r'   r'   r  r(   r     s   
 
;
3	5
4'"5$	N /2&4
(r   c                    s   d gt  j }j D ]\}}	 j|}
t|	||
< qtg |jj}t	 \}}ddl
m} tdd j} fdd|D }j}|ddd	g||}t||   | d
||||||d}|   |j}||_|S )Nr   )
namedtuplec                 S   s   t | dkS rT   )r   )rp   r'   r'   r(   r     s    zast_to_ttir.<locals>.<lambda>c                    s"   i | ]} j |d   j| qS )r   )rb   r7   r4   rd   r`   r'   r(   r    r8   zast_to_ttir.<locals>.<dictcomp>SpecializationProxyr7   r  T)
r   r  r   r  r   r  r	  r
  r  r  )r   rb   r  r   indexr   r   r7   r   r   collectionsr  filterr   ru  r3   r   rv  r  r  )rd   r`   r  r	  r
  r  r  r   r  rp   rf   r  r   r  r  leavesr7   r  proxyr   rF   r'   r  r(   ast_to_ttir  s&   
r  rK   )Gr   r   rr  r"   r  r1  r  dataclassesr   ry   r   typingr   r   r   r   r   r	   r
   r   r   r   r   r   _C.libtritonr   r   r   r   r   language.corer   r   r   runtime.jitr   r   runtimer   _utilsr   r   r   errorsr   r    r!   r)   rG   r   rN   rP   rS   rW   r[   rh   rm   rz   r~   rw   r   rH  r%   r  r   rC  r   r   r   r   r  r'   r'   r'   r(   <module>   s^    ,
W:         ?