a
    yd                    @  s  d dl m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
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 d dlmZmZmZmZmZ d dlZd dlZd dlmZ d dlZd dl m!  m"  mZ# dd	l$m%Z% dd
l&m'Z' dd Z(dd Z)dd Z*G dd dZ+G dd dej,Z-G dd de.Z/G dd de.Z0dd Z1dd Z2dd Z3dd  Z4d!d" Z5d#d$ Z6d%d& Z7d'd( Z8dod)d*d*d+d,d-d.Z9d/d*d0d1d2Z:d/d/d3d4d5Z;ej<d*d6d7d8Z=d9d: Z>ed;d<d=ge? e? gd>Z@d?d@ ZAdAdB ZBdCdD ZCdEdF ZDdGdH ZEdIdJ ZFG dKdL dLZGe< dMdN ZHejIdOdP ZJdQdR ZKdSdT ZLdUdV ZMdpdWdXdYdZd[ZNd\d] ZOd^d_ ZPd`da ZQdbZRdcZSeReReSddZTdeZUdfZVeUeUeVddZWdgdh ZXG didj djZYG dkdl dleZZ[dmdn Z\da]dS )q    )annotationsN)
namedtuple)Path)	get_paths)AnyCallableDictTupleUnion)FileLock   )impl)extractc                 C  s   | d dkr(t | dd  }tj|S tjjtjjtjjtjjtjjtjj	tjj
tjjtjjtjjtjjtjjtjjtjjtjj	d}||  S )Nr   *r   )fp8fp16bf16fp32fp64i1i8i16i32i64u8u16u32u64B)	str_to_tytritonlanguageZpointer_typeZfloat8float16Zbfloat16float32float64int1Zint8int16int32int64Zuint8Zuint16Zuint32Zuint64)nametytys r,   X/var/www/html/stable-diffusion-webui/venv/lib/python3.9/site-packages/triton/compiler.pyr       s(    r   c                 C  s   |   rdt| j S |  r,dt| j S |  r8dS |  rDdS |  rPdS | 	 r\dS | 
 rhdS |  rt| j}dtt| j}| d	| d	S |  rd
S dsJ dd S )NPir   r   r   r   r   _SVFzUnsupported type)Zis_ptr	mangle_tyZ
element_tyZis_intstrZint_bitwidthZis_fp8Zis_fp16Zis_bf16Zis_fp32Zis_fp64Zis_blockZscalarjoinmapshapeZis_void)r*   eltr7   r,   r,   r-   r3   8   s*    
r3   c                   s`   d dd |D }d  fddt D }|dd}|dd}|  d	| d	| }|S )
Nr0   c                 S  s   g | ]}t |qS r,   )r3   .0r*   r,   r,   r-   
<listcomp>R       zmangle_fn.<locals>.<listcomp>c                   s"   g | ]}| d t  |  qS )c)reprr:   r/   	constantsr,   r-   r;   S   r<   .Z_d_'Z_sq___)r5   sortedreplace)r)   Zarg_tysrA   Zmangled_arg_namesZmangled_constantsretr,   r@   r-   	mangle_fnP   s    rH   c                   @  s*   e Zd ZddddZdd Zdd Zd	S )
enter_sub_regionCodeGenerator	generatorc                 C  s
   || _ d S NrK   )selfrL   r,   r,   r-   __init__[   s    zenter_sub_region.__init__c                 C  sL   | j j | _| j j | _i | j _| j j | _| j j	 | _
| j| jfS rM   )rL   lscopecopyliveins
local_defs	prev_defsbuilderget_insertion_blockinsert_blockget_insertion_pointinsert_pointrN   r,   r,   r-   	__enter__^   s    zenter_sub_region.__enter__c                 O  s(   | j j| j | j| j _| j| j _d S rM   )rL   rU   restore_insertion_pointrY   rR   rP   rT   rS   )rN   argskwargsr,   r,   r-   __exit__g   s    
zenter_sub_region.__exit__N)__name__
__module____qualname__rO   r[   r_   r,   r,   r,   r-   rI   Z   s   	rI   c                      s  e Zd Zdde fddZdd Zddd	d
ddZdd Zdd Zdd Z	dd Z
dd Zdd Zdd Zdd Zdd Zdd  Zd!d" Zd#d$ Zd%d& Zd'd( Zd)d* Zd+d, Zd-d. Zd/d0 Zd1d2 Zd3d4 Zd5d6 Zd7d8 Zd9d: Zd;d< Zd=d> Zd?d@ Z dAdB Z!dCdD Z"dEdF Z#dGdH Z$dIdJ Z%dKdL Z&dMdN Z'dOdPdQdRZ(e)j*dSk rndTdU Z+dVdW Z,dXdY Z-dZd[ Z.d\d] Z/d^d_ Z0 fd`daZ1dbdc Z2  Z3S )drJ   NFc
           
      C  s   t j|| _|d u r | j n|| _|	| _|| _|| _t | _	|| _
|| _|| _|| _d | _ttjjtttttd| _g | _i | _i | _d S )N)rangeminfloatintprint
isinstancegetattr)_tritonirrU   create_modulemodulefunction_ret_types	prototypegscopedictrP   
attributesrA   function_name	is_kernel	last_noderc   r    r!   Zminimumre   rf   rg   rh   ri   builtins	scf_stackrS   global_uses)
rN   contextro   rp   rr   rA   rs   rm   rt   function_typesr,   r,   r-   rO   n   s,    	zCodeGenerator.__init__c                 C  sl   d}|| j v r.| j | }|| jvrh|| j|< n:|| jv rD| j| }n$|| jv rZ| j| }nt| d|S )z This function:
        1. make sure `name` is defined
        2. if `name` is triton.language.tensor, get stored tensor by calling
           `self._get_tensor()`
        Nz is not defined)rP   rS   rx   rp   rv   
ValueError)rN   r)   rG   r,   r,   r-   	get_value   s    




zCodeGenerator.get_valuer4   z8Union[triton.language.tensor, triton.language.constexpr]None)r)   valuereturnc                 C  s   || j |< || j|< dS )z This function:
          called by visit_Assign() & visit_FuncDef() to store left value (lvalue)
        1. record local defined name (FIXME: should consider control flow)
        2. store tensor in self.lvalue
        N)rP   rS   )rN   r)   r~   r,   r,   r-   	set_value   s    
zCodeGenerator.set_valuec                 C  s   t |tjjS rM   )rh   r    r!   tensor)rN   r~   r,   r,   r-   is_triton_tensor   s    zCodeGenerator.is_triton_tensorc                 C  s6   |D ] }|  || _t|tjr q&q|o4t|tjS rM   )visitlast_ret_typerh   astReturn)rN   Zstmtsstmtr,   r,   r-   visit_compound_statement   s
    z&CodeGenerator.visit_compound_statementc                 C  s   t j| | d S rM   r   NodeVisitorgeneric_visitrN   noder,   r,   r-   visit_Module   s    zCodeGenerator.visit_Modulec                   s0     |j}|d u sJ  fdd|jD }|S )Nc                   s   g | ]}  |qS r,   r   )r:   r8   rZ   r,   r-   r;      r<   z,CodeGenerator.visit_List.<locals>.<listcomp>)r   ctxelts)rN   r   r   r   r,   rZ   r-   
visit_List   s    zCodeGenerator.visit_Listc                   s     |j}|d u r& jg  d }nrt|trp fdd|D }dd |D } jdd |D  t|}n(tjj	| j} j|j
g |j}|S )Nc                   s   g | ]}t jj| jqS r,   )r    r!   core
_to_tensorrU   r:   vrZ   r,   r-   r;      r<   z.CodeGenerator.visit_Return.<locals>.<listcomp>c                 S  s   g | ]
}|j qS r,   typer   r,   r,   r-   r;      r<   c                 S  s   g | ]
}|j qS r,   handler   r,   r,   r-   r;      r<   )r   r~   rU   rG   rh   tupler    r!   r   r   r   r   )rN   r   Z	ret_valueZret_tyZ
ret_values	ret_typesrG   r,   rZ   r-   visit_Return   s    

zCodeGenerator.visit_Returnc                 C  s:  |  |j\}}t|jjD ]l\}}|jj| d  }|j}|j}tj|t d}	|d u rntj	|	g|d}
ntj
|	||d}
|  |
 q| jrdnd}| j| j| j| j| j|}| j| | }g }d}t|D ]\}}|| jv r*| j| }t|tjjstj| j| }|| qq|| jv rN||d| j| d  |tj||| jj|  |d7 }q| j }t||D ]\}}|  || q| j!| | "|j#}|s| j$g  nRt| j%t&rt'| j%| j_(|)| j| j n | j%g| j_(|)| j| j |r6| j*| d S )	Nr   idr   targetsr~   )targetr~   
annotationpublicprivater   ztt.divisibility)+r   r]   	enumeratedefaultsr   argr   NameStoreAssign	AnnAssignrt   rU   Zget_or_insert_functionrm   rs   ro   to_irZ	push_backZadd_entry_blockrA   rh   r    r!   	constexprappendrr   Zset_arg_attrr   Zparam_typesrV   zipr   set_insertion_point_to_startr   bodyrG   r   r   listr   Z
reset_typeset_insertion_point_to_end)rN   r   	arg_nameskwarg_namesr/   default_valueZarg_noder   r)   Z	st_targetZ	init_nodeZ
visibilityfnentryZ
arg_valuesidxZarg_nameZcstZ	insert_pt	arg_valueZhas_retr,   r,   r-   visit_FunctionDef   sR     

"

zCodeGenerator.visit_FunctionDefc                 C  s4   g }|j D ]}|| |g7 }q
| |j}||fS rM   )r]   r   kwarg)rN   r   r   r   r   r,   r,   r-   visit_arguments  s
    
zCodeGenerator.visit_argumentsc                 C  s   t j| | |jS rM   )r   r   r   r   r   r,   r,   r-   	visit_arg  s    zCodeGenerator.visit_argc                 C  s   |  |j}|  |j}|  |j}|tjjkrv|| jv rHt| dt	|tjjsbtj|}|| j|< | j| S | 
|S )Nz4 is already defined. constexpr cannot be reassigned.)r   r   r   r~   r    r!   r   rP   r{   rh   visit_Assign)rN   r   r   r   r~   r,   r,   r-   visit_AnnAssign  s    


zCodeGenerator.visit_AnnAssignc                 C  s   g }|j D ]}|| |g7 }q
t|dks0J |d }| |j}t|tsT|g}t|tsd|g}t||D ]H\}}t|tjj	r|j}t|tjj
stjj|| j}| || qnd S )Nr   r   )r   r   lenr~   rh   r   r   r    r!   r   r   r   r   rU   r   )rN   r   Z_namesr   namesvaluesr)   r~   r,   r,   r-   r   '  s     


zCodeGenerator.visit_Assignc                 C  sR   |j j}tj|t d}t||j|j}tj|j g|d}| 	| | 
|S )Nr   r   )r   r   r   r   LoadBinOpopr~   r   r   r|   )rN   r   r)   lhsrhsZassignr,   r,   r-   visit_AugAssign:  s    
zCodeGenerator.visit_AugAssignc                 C  s"   t |jtjkr|jS | |jS rM   )r   r   r   r   r   r|   r   r,   r,   r-   
visit_NameB  s    zCodeGenerator.visit_Namec                 C  s   t j| | d S rM   r   r   r,   r,   r-   visit_StoreG  s    zCodeGenerator.visit_Storec                 C  s   t j| | d S rM   r   r   r,   r,   r-   
visit_LoadJ  s    zCodeGenerator.visit_Loadc                   s    fdd|j D }t|S )Nc                   s   g | ]}  |qS r,   r   )r:   xrZ   r,   r-   r;   N  r<   z-CodeGenerator.visit_Tuple.<locals>.<listcomp>)r   r   )rN   r   r]   r,   rZ   r-   visit_TupleM  s    zCodeGenerator.visit_Tuplec                 C  s   |  |j}|  |j}tjdtjdtjdtjdtjdtj	dtj
dtjdtjd	tjd
tjdtjdit|j }| |rt|||| jdS | |r|d d d |dd   }t|||| jdS t|||S d S )N__add____sub____mul____truediv____floordiv____mod____pow__
__lshift__
__rshift____and____or____xor___builder   r)r   leftrightr   AddSubMultDivFloorDivModPowLShiftRShiftBitAndBitOrBitXorr   r   r   ri   rU   rN   r   r   r   r   r,   r,   r-   visit_BinOpQ  s.    

zCodeGenerator.visit_BinOpc                 C  s.  | j | | |j | j  }| j }i }|jrr| j | | | _i | _| |j | j }| j  }g }g }g }	|D ]}
|df|dffD ]P\}}|
|v r||
 j	||
 j	ksJ d|
 d||
 j	 d| d||
 j	 q|
|v s|
|v rV|
|
 |
|
|v r||
 j	n||
 j	 |	
|
|v rF||
 j n||
 j  |
|v rv|
|vrv||
 ||
< |
|v r|
|vr||
 ||
< q| | @ D ]v}
|
|v rq||
 j	}||
 j	}||ksJ d|
 d| d	| d
|
|
 |
| |	
||
 j  q|||||||	fS )NZ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 ())rU   r   r   r   rV   rS   rQ   orelserP   r   r   r   get_typekeys)rN   r   rR   
then_block
else_block	then_defs	else_defsr   r   ir_ret_typesr)   ZdefsZ
block_nameZthen_tyZelse_tyr,   r,   r-   visit_then_else_blocksj  s^    






$,




z$CodeGenerator.visit_then_else_blocksc                   sH  t | }|\}}| j }| j }| j }| j| | j|j|| | ||||\ }}}	}
}| j| | s| j|fdd|	D  | j| | s| j| fdd|	D  |D ]}|	| qW d    n1 s0    Y  | j
| t|	D ]0\}}tjj|||
| }| || qd S )Nc                   s   g | ]} | j qS r,   r   r:   nr   r,   r-   r;     r<   z4CodeGenerator.visit_if_top_level.<locals>.<listcomp>c                   s   g | ]} | j qS r,   r   r   r   r,   r-   r;     r<   )rI   rU   create_blockr   Zcreate_cond_branchr   r   Zhas_terminatorZcreate_branchadd_argumentr   r   r    r!   r   r   r   r   )rN   condr   srrR   Zip_blockr   r   Zendif_blockr   r   r   r*   r/   r)   
new_tensorr,   )r   r   r-   visit_if_top_level  s*    



*z CodeGenerator.visit_if_top_levelc                   s~  t $}|\}}j }j }|jr8j nd }||||\ }}}	}
}j| jfdd|
D |jd}|	|
  j|
  t|	dkrʈjfdd|	D  |js| }n|	|  j|  t|	dkr j fdd|	D  W d    n1 s60    Y  t|	D ]0\}}tjj|||
| }|| qHd S )Nc                   s   g | ]}|  jqS r,   r   rU   r9   rZ   r,   r-   r;     r<   z.CodeGenerator.visit_if_scf.<locals>.<listcomp>Tr   c                   s   g | ]} | j qS r,   r   r   r   r,   r-   r;     r<   c                   s   g | ]} | j qS r,   r   r   r   r,   r-   r;     r<   )rI   rU   rX   r   r   r   r\   Zcreate_if_opr   Zmerge_block_beforeZget_then_blockr   r   create_yield_opZget_else_blockr   r    r!   r   r   
get_resultr   )rN   r  r   r  rR   r0   ipr   r   r   r   Zif_opr/   r)   r  r,   )r   rN   r   r-   visit_if_scf  s,    

 
:zCodeGenerator.visit_if_scfc                 C  s   |  |j}t|tjjrP|jtjj| jd}| j	rB| 
|| q| || n2t|tjjrd|j}|rv| |j n| |j d S )Nr   )r   testrh   r    r!   r   tor%   rU   rw   r	  r  r   r~   r   r   r   rN   r   r  r,   r,   r-   visit_If  s    zCodeGenerator.visit_Ifc                 C  s.   |  |j}|jr|  |jS |  |jS d S rM   )r   r
  r~   r   r   r  r,   r,   r-   visit_IfExp  s    zCodeGenerator.visit_IfExpc                 C  s   d S rM   r,   r   r,   r,   r-   
visit_Pass  s    zCodeGenerator.visit_Passc                 C  sT  t |jdksJ t |jdks$J | |j}| |jd }t|tjjrT|j	}t|tjjrh|j	}t
|jd tjkrtj||u S t
|jd tjkrtj||uS tjdtjdtjdtjdtjdtjdit
|jd  }| |rt|||| jd	S | |rB|d d
 d |d
d   }t|||| jd	S t|||S d S )Nr   r   __eq____ne____lt____le____gt____ge__r   r   r   )r   comparatorsopsr   r   rh   r    r!   r   r~   r   r   IsIsNotEqNotEqLtLtEGtGtEr   ri   rU   r   r,   r,   r-   visit_Compare  s6    zCodeGenerator.visit_Comparec                 C  sZ   |  |j}tjdtjdtjdtjdit|j }| 	|rNt
||| jdS t
|| S )N__neg____pos____not__
__invert__r   )r   operandr   USubUAddNotInvertr   r   r   ri   rU   )rN   r   r   r   r,   r,   r-   visit_UnaryOp  s    
zCodeGenerator.visit_UnaryOpc                   sF  t }|\}}j }j| j| |j j  j	}g }g }g }	|D ]p}
|
|v rb
||
 sJ 
||
 sJ ||
 j||
 jksJ ||
 |||
 j |	||
  qbj| jfdd|D dd |	D }j| fdd|D  j  t|D ]:\}}
tjj ||| j|
< j|
 j	|
< q:|j}j  j|j fddtt|	D  j| fdd|D }j| t|D ]:\}}
tjj|||| j|
< j|
 j	|
< qj| |j j  j	}g }|D ]}
|
|v rT|||
  qTjdd |D  W d    n1 s0    Y  t|D ]"\}}
||	| j|| qt|D ]8\}}
tjj| ||| }|j|
< |j	|
< q|j!D ]"}ds0J d	t"j#$| qd S )
Nc                   s   g | ]}|  jqS r,   r  r9   rZ   r,   r-   r;   :  r<   z-CodeGenerator.visit_While.<locals>.<listcomp>c                 S  s   g | ]
}|j qS r,   r   r:   r   r,   r,   r-   r;   ;  r<   c                   s   g | ]}|  jqS r,   r  r9   rZ   r,   r-   r;   >  r<   c                   s   g | ]}  |qS r,   )r   r?   )before_blockr,   r-   r;   F  r<   c                   s   g | ]}|  jqS r,   r  r9   rZ   r,   r-   r;   I  r<   c                 S  s   g | ]
}|j qS r,   r   r:   yr,   r,   r-   r;   X  r<   FzNot implemented)%rI   rU   r   r   rw   r   r   r   poprS   r   r   r   Zcreate_while_opZcreate_block_with_parentZ
get_beforer   r    r!   r   r   r   rP   r   r
  Zcreate_condition_opr   rc   r   Z	get_afterr  Zreplace_use_in_block_withr  r   r   r   r   )rN   r   r  rR   rW   dummyZ	loop_defsr   r   	init_argsr)   Zwhile_opr/   r  Zafter_blockyieldsZnew_defr   r,   )r,  rN   r-   visit_While  st    


 & 

6

zCodeGenerator.visit_Whilec                 C  sL   |j jjdksJ | |j}| |j}| |rD|j|| jdS || S )Nr   r   )	r   	__class__r`   r   r~   slicer   __getitem__rU   )rN   r   r   Zslicesr,   r,   r-   visit_Subscripth  s    
zCodeGenerator.visit_Subscriptc                   s    fdd|j D S )Nc                   s   g | ]}  |qS r,   r   )r:   ZdimrZ   r,   r-   r;   q  r<   z0CodeGenerator.visit_ExtSlice.<locals>.<listcomp>)dimsr   r,   rZ   r-   visit_ExtSlicep  s    zCodeGenerator.visit_ExtSlicec                   sN    |jj} fdd|jjD }|tjjkr|| }t|jj	|j
j	|jj	}|D ]@}tj| j|jj<  |j |jD ]}tj | qqTd S | jd krtdt|dkr|d n  td}t|dkr|d n  |jjd }	t|dkr|d n  td}
d}t|
tjjr^|
j	dk r^tj|
j	 }
d	}|	| }}	tjj| jj}tjj|	 jj}	tjj|
 jj}
 j|} j|	}	 j|
}
 j  j! } "|jjtjj#|tjjj$ t% }|\}} j& } j' } j(|  j)*|  |j  j)+  |,  g }g }g } j-D ]}||v rb . j-| sJ | d
 .|| sJ  j-| j/|| j/ksJ d| d|| j/ d j-| j/ d|*| |*tjj||  j |*tjj j-|  j qb j0|  j1||	|
dd |D } j)*|  j(|2d t3|D ]8\}} "|tjj#|2d4|d || j/ q| |j  j)+  g } j-D ].}||v r|*tjj j-|  j qt|dkr* j5dd |D  |2d6 }|7 dksNJ d j(|2d  j8|9 }|r j8|	} j:||} j|jj j;|  "|jjtjj#|tjjj$ W d    n1 s0    Y  t3|D ].\}} "|tjj#|<||| j/ q|jD ]"}ds8J dtj | q&d S )Nc                   s   g | ]}  |qS r,   r   r+  rZ   r,   r-   r;   u  r<   z+CodeGenerator.visit_For.<locals>.<listcomp>rc   zAOnly `range` and `static_range` iterators are currently supportedr   r   r   FTz is not tensorzLoop-carried variable z has initial type z but is re-assigned to z: in loop! Please make sure that the type stays consistent.c                 S  s   g | ]
}|j qS r,   r   r+  r,   r,   r-   r;     r<   c                 S  s   g | ]
}|j qS r,   r   r-  r,   r,   r-   r;     r<   z7We use SCF, so the loop body should only have one blockz)Don't know what to do with else after for)=r   iterfuncr]   r    r!   static_rangerc   startr~   endstepr   rP   r   r   r   r   r   r   r   r   rv   RuntimeErrorr   Numrh   r   r   rU   r   Zcreate_to_indexZcreate_undefZget_int32_tyr   r   r'   rI   rX   r   r   rw   r   r/  ZeraserS   r   r   r\   Zcreate_for_opget_bodyr   r   r  Z
get_parentsizeZcreate_index_to_siZget_induction_varZ
create_subZreplace_all_uses_withr  )rN   r   ZIteratorClassZ	iter_argsiteratorr<  r/   r   ZlbZubr?  Znegative_stepZivr  rR   rW   r  blockr1  r2  r   r)   Zfor_opZfor_op_regionZub_sir,   rZ   r-   	visit_Fors  s    
$&&
"




 

"2


"B(
zCodeGenerator.visit_Forc                 C  s0   |  |j}|  |j}|  |j}t|||S rM   )r   lowerupperr?  r5  )rN   r   rG  rH  r?  r,   r,   r-   visit_Slice  s    zCodeGenerator.visit_Slicec                 C  s   |  |jS rM   )r   r~   r   r,   r,   r-   visit_Index  s    zCodeGenerator.visit_Indexc                 C  s   |j | |jiS rM   )r   r   r~   r   r,   r,   r-   visit_keyword  s    zCodeGenerator.visit_keywordc              
     s   |j}t|tjjr |j}t }|jD ]}|	 | q,fdd|j
D  t|tjjr"ddlm} ||jg R i |  fdd|jD  dd  D  t }dd t D  fdd	D }fd
dt D  dd  D }dd  D }	t|j|	|}
j|
stjg |	}tj|jj j}tjj||||j|
jd}| |  |j }|j|
< n
j|
 }j!|
}j"||}|# dks|d u rd S |# dkrtj$|%d|S g }t&|# D ]$}|'tj$|%|||  qt(|S t)|dr<*|j+sHt,-|r\| dji|S |j./ v rzdd  D  | i |S )Nc                   s   g | ]}  |qS r,   r   r+  rZ   r,   r-   r;     r<   z,CodeGenerator.visit_Call.<locals>.<listcomp>r   )getcallargsc                   s   g | ]} | qS r,   r,   )r:   r)   r]   r,   r-   r;     r<   c                 S  s*   g | ]"}t |tjjr|n
tj|qS r,   )rh   r    r!   r   r   r+  r,   r,   r-   r;     s   c                 S  s"   g | ]\}}t |tjjr|qS r,   )rh   r    r!   r   r:   r/   r   r,   r,   r-   r;     r<   c                   s   i | ]}| | qS r,   r,   r?   rM  r,   r-   
<dictcomp>  r<   z,CodeGenerator.visit_Call.<locals>.<dictcomp>c                   s    g | ]\}}| v rd n|qS rM   r,   rN  )
constexprsr,   r-   r;     r<   c                 S  s   g | ]}|d ur|j qS rM   r   r+  r,   r,   r-   r;     r<   c                 S  s   g | ]}|d ur|j qS rM   r   r+  r,   r,   r-   r;     r<   )rm   rs   rz   r   __self__r   c                 S  s$   g | ]}t |tjjr|jn|qS r,   )rh   r    r!   r   r~   r+  r,   r,   r-   r;     s   )0r   r;  rh   r    r!   r   r~   rq   keywordsupdater]   runtimeJITFunctioninspectrL  r   r   r   rH   r`   rm   Zhas_functionfunction_typesysmodulesra   __dict__rJ   rU   ry   rn   parser   Zget_functioncallZget_num_resultsr   r  rc   r   r   hasattrr   rQ  r   
is_builtinrv   r   )rN   r   r   kwskeywordrL  rr   rA   arg_vals	arg_typesfn_namero   rp   rL   Zcallee_ret_typesymbolZcall_opresultsr/   r,   )r]   rP  rN   r-   
visit_Call  s`    
 
"zCodeGenerator.visit_Callc                 C  s   t j|jS rM   r    r!   r   r~   r   r,   r,   r-   visit_Constant#  s    zCodeGenerator.visit_Constantz
ast.BoolOp)r   c                 C  s   t |jdksJ | |jd }| |jd }tjdtjdit|j }| |rjt	|||| j
dS | |r|d d d |dd   }t	|||| j
dS t	|||S d S )Nr   r   r   logical_and
logical_orr   r   )r   r   r   r   AndOrr   r   r   ri   rU   r   r,   r,   r-   visit_BoolOp&  s    

zCodeGenerator.visit_BoolOp)      c                 C  s   t j|jS rM   rg  r   r,   r,   r-   visit_NameConstant9  s    z CodeGenerator.visit_NameConstantc                 C  s   t j|jS rM   )r    r!   r   r   r   r,   r,   r-   	visit_Num<  s    zCodeGenerator.visit_Numc                 C  s   t jt|S rM   )r    r!   r   r   literal_evalr   r,   r,   r-   	visit_Str?  s    zCodeGenerator.visit_Strc                 C  sD   |  |j}t|tjjr8|jdkr8tjjj|| j	dS t
||jS )NT)rU   )r   r~   rh   r    r!   r   attrZsemanticZtransrU   ri   )rN   r   r   r,   r,   r-   visit_AttributeB  s
    
zCodeGenerator.visit_Attributec                 C  s   t j| | d S rM   r   r   r,   r,   r-   
visit_ExprI  s    zCodeGenerator.visit_Exprc                 C  s   d S rM   r,   r   r,   r,   r-   visit_NoneTypeL  s    zCodeGenerator.visit_NoneTypec                   s^   |d ur|| _ t 4 tdt tdt t |W  d    S 1 sP0    Y  d S )Nignore)ru   warningscatch_warningssimplefilterDeprecationWarningPendingDeprecationWarningsuperr   r   r4  r,   r-   r   O  s    
zCodeGenerator.visitc                 C  s   t |j}td|d S )NzUnsupported node: {})r   r`   NotImplementedErrorformat)rN   r   typenamer,   r,   r-   r   Y  s    
zCodeGenerator.generic_visit)4r`   ra   rb   rq   rO   r|   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  r	  r  r  r  r   r*  r3  r7  r9  rF  rI  rJ  rK  rf  rh  rm  rX  version_inforp  rq  rs  rv  rw  rx  r   r   __classcell__r,   r,   r  r-   rJ   m   s\   
56Ko5
rJ   c                      s$   e Zd Z fddZdd Z  ZS )CompilationErrorc                   sv   d|j  d|j d| _|  jd|dd |j  7  _|  jdd|j  d 7  _|| _|| _t | j d S )Nzat :z:

 ^)	lineno
col_offsetmessager5   splitsrcr   r  rO   )rN   r  r   r  r,   r-   rO   _  s    $zCompilationError.__init__c                 C  s   t | | j| jffS rM   )r   r  r   rZ   r,   r,   r-   
__reduce__g  s    zCompilationError.__reduce__r`   ra   rb   rO   r  r  r,   r,   r  r-   r  ^  s   r  c                      s$   e Zd Z fddZdd Z  ZS )OutOfResourcesc                   sJ   d| d| d| | _ |  j d7  _ || _|| _|| _t | j  d S )Nzout of resource: z, Required: z, Hardware limit: z0. Reducing block sizes or `num_stages` may help.)r  requiredlimitr)   r  rO   )rN   r  r  r)   r  r,   r-   rO   m  s    zOutOfResources.__init__c                 C  s   t | | j| j| jffS rM   )r   r  r  r)   rZ   r,   r,   r-   r  w  s    zOutOfResources.__reduce__r  r,   r,   r  r-   r  l  s   
r  c                 C  sJ   d}t | D ]8\}}|t|7 }||jv r2|d7 }||jv r|d7 }q|S )N r=   d)r   r4   
equal_to_1divisible_by_16)	signaturespecializationsuffixr/   r0   r,   r,   r-   kernel_suffix|  s    


r  c                 C  s   t j| |}||_|S rM   )rj   rk   parse_mlir_modulery   )pathry   rm   r,   r,   r-   r    s    r  c              
     sv  t |tr"dd t|dD }tj }|  fddfdd  D  j	
 }djt| |g}t| fdd|jD }d	d |jD } 
 }	|	|  fd
d| D }
tjg |
}t||||	||dd}z|  W nX ty` } z>|j}|d u s:t |ttfr>|tj||W Y d }~n
d }~0 0 |j}||_||fS )Nc                 S  s   i | ]\}}||  qS r,   stripr:   kr   r,   r,   r-   rO    r<   z#build_triton_ir.<locals>.<dictcomp>,c                   s   t | tr j| S | S rM   )rh   r4   r   index)r/   r   r,   r-   <lambda>  r<   z!build_triton_ir.<locals>.<lambda>c                   s   i | ]\}} ||qS r,   r,   )r:   keyr~   )cst_keyr,   r-   rO    r<   r0   c                   s*   i | ]"}|| v r" | d kr"dndqS )r   Tr   r,   r:   r  )r+   r,   r-   rO    r<   c                 S  s   i | ]
}|d qS ))Zmultiple_of   r,   r  r,   r,   r-   rO    r<   c                   s    g | ]\}}| vrt |qS r,   )r   r  r@   r,   r-   r;     r<   z#build_triton_ir.<locals>.<listcomp>T)rp   rA   rs   rr   rt   )rh   r4   r   r  rj   rk   ry   Zload_tritonitems__globals__rQ   r5   r`   r  r   r   r  r  rS  r    r!   rW  rJ   r   r[  	Exceptionru   r  r  r  rm   )r   r  r  rA   ry   rp   rs   Znew_constantsZ	new_attrsZall_constantsrb  ro   rL   er   rG   r,   )rA   r  r   r+   r-   build_triton_ir  s4    



$r  c                 C  sL   t j| j}|  |  |  |  |  |	  |
|  | S rM   )rj   rk   pass_managerry   enable_debugZadd_inliner_passZadd_triton_combine_passadd_canonicalizer_passadd_cse_passadd_licm_passrun)modpmr,   r,   r-   optimize_triton_ir  s    
r  c                 C  s   t | |||\}}t|S rM   )r  r  )r   r  r  rA   r  r0   r,   r,   r-   ast_to_ttir  s    r  c                 C  s   t j| j}|| |  |  || || |	  |
  |  || |  || |  |  |d dkr|  |  |  |  ||  | S )N
      )rj   rk   r  ry   Z$add_convert_triton_to_tritongpu_passr  Zadd_coalesce_passZadd_tritongpu_combine_passZadd_tritongpu_pipeline_passZadd_tritongpu_prefetch_passr  r  r  Z(add_tritongpu_decompose_conversions_passZ'add_tritongpu_update_mma_for_volta_passZadd_symbol_dce_passZ'add_tritongpu_reorder_instructions_passr  )r  	num_warps
num_stagescompute_capabilityr  r,   r,   r-   ttir_to_ttgir  s*    





r  c                 C  sR   |  D ]&\}}t|dks(t|dkr d S qt| t| t|  d S Nr   )r  r   rj   add_external_libsr   r   r   )r  Zlibsr)   r  r,   r,   r-   r    s    r  c                 C  s   |rt | | t| |S rM   )r  rj   Ztranslate_triton_gpu_to_llvmir)r  extern_libsr  r,   r,   r-   ttgir_to_llir  s    
r  r   rf   zTuple[str, int])r  r  ptx_versionr   c                 C  s(   |du rt  \}}t|}t| ||S )z
    Translate TritonGPU module to PTX code.
    :param mod: a TritonGPU dialect module
    :return:
        - PTX code
        - shared memory allocation size
    N)path_to_ptxasptx_get_versionrj   Ztranslate_llvmir_to_ptx)r  r  r  r0   cuda_versionr,   r,   r-   llir_to_ptx  s    
r  r4   )ptxr  c                 C  s   t  \}}t| ||S )z
    Compile TritonGPU module to cubin.
    :param ptx: ptx code
    :param compute_capability: compute capability
    :return: str
    )r  rj   Zcompile_ptx_to_cubin)r  r  ptxasr0   r,   r,   r-   ptx_to_cubin  s    
r  )r  r   c                 C  s>   | sJ |  dD ]&}| }|dr|  d   S qdS )zd
    Get kernel name from PTX code.
    This Kernel name is required when launching the kernel.
    r  z	// .globlN)r  r  
startswith)r  liner,   r,   r-   ptx_get_kernel_name
  s
    
r  )r   c                 C  s^   t | tsJ tt| d\}}|dkr2d| S |dkrBd| S |dkrRd| S tdd	S )
zK
    Get the highest PTX version supported by the current CUDA driver.
    rB      P      F   r  ?   z'Triton only support CUDA 10.0 or higherN)rh   r4   r6   rf   r  r@  )r  majorminorr,   r,   r-   r    s    r  c                  C  s   t jt} t jddt j| ddddg}|D ]j}t j|r2t j|r2t	j
|dgt	jd}|d ur2tjd	|d
tjd}|d ur2||df  S q2tdd S )NZTRITON_PTXAS_PATHr  Zthird_partycudabinr  z	--version)stderrz.*release (\d+\.\d+).*utf-8)flagsr   zCannot find ptxas)osr  dirname__file__environgetr5   existsisfile
subprocesscheck_outputSTDOUTresearchdecode	MULTILINEgroupr@  )base_dirpathsr  resultversionr,   r,   r-   r  '  s    r  instance_descriptorr  r  )r   c                 C  s2   | d dkrdS ddddddd	d
d
d
d
dd|  S )Nr   r   ZCUdeviceptrint32_tZint8_tZint16_tint64_tuint32_tuint64_tre   double)r   r   r   r   r   r   r   r   r   r   f32r   r,   r*   r,   r,   r-   	ty_to_cpp@  s"    r  c                 C  s(   d}|  d}t|D ]\}}| qd S )Nzint i = 0;
r  )r  r   )r  r  r+   r/   r*   r,   r,   r-   generate_name_initializerS  s    
r  c                 C  s.   t | dkr$dt| d  } |  dS )N   Zkernel_r  z.h)r   hashlibsha256encode	hexdigest)r)   r,   r,   r-   binary_name_to_header_nameZ  s    r  c                   s   d dd | D }dd  dd dd	  fd
d| D  }d| dd fdd| D  dd  fdd| D  d| dd dd | D  dd dd | D  dd dd | D  d}|S )N, c                 s  s$   | ]\}}t | d | V  qdS )z argN)r  r:   r/   r*   r,   r,   r-   	<genexpr>b  r<   z$generate_launcher.<locals>.<genexpr>c                 S  s.   | d dkrdS dddddddddd	d

|  S )Nr   r   	PyObject*r  r  r  r  re   r  )
r   r   r   r   r   r   r   r   r  r   r,   r  r,   r,   r-   _extracted_typed  s    z*generate_launcher.<locals>._extracted_typec              	   S  s   ddddddddd	|  S )
NOfr  lIr/   KL)r  re   r  longr  r  r  r  r,   r  r,   r,   r-   	format_oft  s    	z$generate_launcher.<locals>.format_ofZ
iiiiiKKOOOr  c                   s   g | ]} |qS r,   r,   r9   )r  r  r,   r-   r;     r<   z%generate_launcher.<locals>.<listcomp>aX  
#include "cuda.h"
#include <stdbool.h>
#include <Python.h>

static inline void gpuAssert(CUresult code, const char *file, int line)
{
   if (code != CUDA_SUCCESS)
   {
      const char* prefix = "Triton Error [CUDA]: ";
      const char* str;
      cuGetErrorString(code, &str);
      char err[1024] = {0};
      strcat(err, prefix);
      strcat(err, str);
      PyErr_SetString(PyExc_RuntimeError, err);
   }
}

#define CUDA_CHECK(ans) { gpuAssert((ans), __FILE__, __LINE__); }

void _launch(int gridX, int gridY, int gridZ, int num_warps, int shared_memory, CUstream stream, CUfunction function, z) {
  void *params[] = { c                 3  s    | ]}| vrd | V  qdS )z&argNr,   r?   r@   r,   r-   r    r<   a   };
  if(gridX*gridY*gridZ > 0){
    CUDA_CHECK(cuLaunchKernel(function, gridX, gridY, gridZ, 32*num_warps, 1, 1, shared_memory, stream, params, 0));
  }
}

typedef struct _DevicePtrInfo {
    CUdeviceptr dev_ptr;
    bool valid;
} DevicePtrInfo;

static inline DevicePtrInfo getPointer(PyObject *obj, int idx) {
  DevicePtrInfo ptr_info;
  ptr_info.dev_ptr = 0;
  ptr_info.valid = true;
  if (PyLong_Check(obj)) {
    ptr_info.dev_ptr = PyLong_AsUnsignedLongLong(obj);
    return ptr_info;
  }
  if (obj == Py_None) {
    // valid nullptr
    return ptr_info;
  }
  PyObject *ptr = PyObject_GetAttrString(obj, "data_ptr");
  if(ptr){
    PyObject *empty_tuple = PyTuple_New(0);
    PyObject *ret = PyObject_Call(ptr, empty_tuple, NULL);
    Py_DECREF(empty_tuple);
    Py_DECREF(ptr);
    if (!PyLong_Check(ret)) {
      PyErr_SetString(PyExc_TypeError, "data_ptr method of Pointer object must return 64-bit int");
      ptr_info.valid = false;
      return ptr_info;
    }
    ptr_info.dev_ptr = PyLong_AsUnsignedLongLong(ret);
    unsigned attr;
    CUresult status =
        cuPointerGetAttribute(&attr, CU_POINTER_ATTRIBUTE_MEMORY_TYPE, ptr_info.dev_ptr);
    if (ptr_info.dev_ptr &&
        (!(attr == CU_MEMORYTYPE_DEVICE || attr == CU_MEMORYTYPE_UNIFIED) ||
         !(status == CUDA_SUCCESS))) {
        PyErr_Format(PyExc_ValueError,
                     "Pointer argument (at %d) cannot be accessed from Triton (cpu tensor?)", idx);
        ptr_info.valid = false;
    }
    Py_DECREF(ret);  // Thanks ChatGPT!
    return ptr_info;
  }
  PyErr_SetString(PyExc_TypeError, "Pointer argument must be either uint64 or have data_ptr method");
  return ptr_info;
}

static PyObject* launch(PyObject* self, PyObject* args) {
  int gridX, gridY, gridZ;
  uint64_t _stream;
  uint64_t _function;
  int num_warps;
  int shared_memory;
  PyObject *launch_enter_hook = NULL;
  PyObject *launch_exit_hook = NULL;
  PyObject *compiled_kernel = NULL;
  PyObject *hook_ret = NULL;
  r  c                   s$   g | ]\}} | d | dqS )z _arg; r,   r   )r  r,   r-   r;     r<   z
  if(!PyArg_ParseTuple(args, "z", &gridX, &gridY, &gridZ, &num_warps, &shared_memory, &_stream, &_function, &launch_enter_hook, &launch_exit_hook, &compiled_kernel, c                 s  s   | ]\}}d | V  qdS )z&_argNr,   r   r,   r,   r-   r    r<   z)) {
    return NULL;
  }

  if (launch_enter_hook != Py_None) {
    PyObject *new_args = PyTuple_Pack(1, compiled_kernel);
    hook_ret = PyObject_CallObject(launch_enter_hook, new_args);
    Py_DECREF(new_args);
  }


  // raise exception asap
  r  c                 S  s>   g | ]6\}}|d  dkr6d| d| d| d| d	ndqS )r   r   zDevicePtrInfo ptr_infoz = getPointer(_argr  z); if (!ptr_infoz.valid) return NULL;r  r,   r   r,   r,   r-   r;     r<   ze;
  _launch(gridX, gridY, gridZ, num_warps, shared_memory, (CUstream)_stream, (CUfunction)_function, c                 s  s4   | ],\}}|d  dkr"d| dnd| V  qdS )r   r   Zptr_infoz.dev_ptrZ_argNr,   r   r,   r,   r-   r    r<   a  );

  if (launch_exit_hook != Py_None) {
    PyObject *new_args = NULL;
    if (hook_ret) {
        new_args = PyTuple_Pack(2, compiled_kernel, hook_ret);
    } else {
        new_args = PyTuple_Pack(1, compiled_kernel);
    }
    hook_ret = PyObject_CallObject(launch_exit_hook, new_args);
    Py_DECREF(new_args);
  }

  if (hook_ret) {
      Py_DECREF(hook_ret);
  }
  if(PyErr_Occurred()) {
    return NULL;
  }
  // return None
  Py_INCREF(Py_None);
  return Py_None;
}

static PyMethodDef ModuleMethods[] = {
  {"launch", launch, METH_VARARGS, "Entry point for all kernels with this signature"},
  {NULL, NULL, 0, NULL} // sentinel
};

static struct PyModuleDef ModuleDef = {
  PyModuleDef_HEAD_INIT,
  "__triton_launcher",
  NULL, //documentation
  -1, //size
  ModuleMethods
};

PyMODINIT_FUNC PyInit___triton_launcher(void) {
  PyObject *m = PyModule_Create(&ModuleDef);
  if(m == NULL) {
    return NULL;
  }
  PyModule_AddFunctions(m, ModuleMethods);
  return m;
}
)r5   r  r   r   )rA   r  Z	arg_declsr  r  r,   )r  rA   r  r-   generate_launchera  s*    "TUUab r  c                   C  s   t jt jd ddS )NHOMEz.tritoncache)r  r  r5   r  r,   r,   r,   r-   default_cache_dir  s    r  c                  C  s   d} t jd| dS )Nz/usr/local/cudaZ	CUDA_HOME)default)r  getenv)Zdefault_dirr,   r,   r-   default_cuda_dir  s    r  c                   @  s.   e Zd Zdd Zdd Zdd Zddd	Zd
S )CacheManagerc                 C  s^   || _ d | _tjdt | _| jrZtj| j| j | _tj| jd| _tj	| jdd d S )NZTRITON_CACHE_DIRlockT)exist_ok)
r  	lock_pathr  r  r  r  	cache_dirr  r5   makedirs)rN   r  r,   r,   r-   rO   "  s    zCacheManager.__init__c                 C  s   t j| j|S rM   )r  r  r5   r  rN   filenamer,   r,   r-   
_make_path,  s    zCacheManager._make_pathc                 C  s   | j s
dS tj| |S )NF)r  r  r  r  r  r  r,   r,   r-   has_file/  s    zCacheManager.has_fileTc              	   C  s   | j s
d S t|t}|s t|}| jd us.J | |}t| jd |rLdnd}t|d |}|| W d    n1 s~0    Y  t	
|d | W d    n1 s0    Y  d S )Nwbwz.tmp)r  rh   bytesr4   r  r  r   openwriter  rename)rN   datar  binaryfilepathmoder  r,   r,   r-   put4  s    

(zCacheManager.putN)T)r`   ra   rb   rO   r  r  r(  r,   r,   r,   r-   r     s   
r  c                  C  s0   t ddg   dd  } dd | D S )NZwhereisz
libcuda.sor   c                 S  s   g | ]}t j|qS r,   )r  r  r  )r:   locr,   r,   r-   r;   J  r<   z libcuda_dirs.<locals>.<listcomp>)r  r  r  r  r  )Zlocsr,   r,   r-   libcuda_dirsG  s    "r*  c               	   c  sR   t jt j } }t t  t _t _zd V  W | | t _t _n| | t _t _0 d S rM   )rX  stdoutr  ioStringIO)Z
old_stdoutZ
old_stderrr,   r,   r-   quietM  s
    r.  c                 C  s  t  }tjdt }tj|d}tjt}tj|d}tj|d}tj|d}	tj	|svtj	|	rv|}t
d}
tj|dj| |
d}tjd}|d u rtd	}td
}|d ur|n|}|d u rtdt d }||dd| d| d| dddd|g}|dd |D 7 }t|}|dkr@|S g }|}||g}dg}g }tj| d|g||dg |||d}dg}|d|  |d|  |d t| |g|d}t   tjf i | W d    n1 s0    Y  |S )NZ	CUDA_PATHincludezthird_party/cuda/includezcuda.h
EXT_SUFFIXz{name}{suffix})r)   r  CCclanggcczFFailed to find C compiler. Please specify via CC environment variable.z-O3z-Iz-sharedz-fPICz-lcudaz-oc                 S  s   g | ]}d | qS )z-Lr,   )r:   dirr,   r,   r-   r;   o  r<   z_build.<locals>.<listcomp>r   r  r=   )r)   r!   sourcesinclude_dirsextra_compile_argsextra_link_argslibrary_dirs	libraries	build_extz--build-temp=z--build-lib=z-q)r)   ext_modulesscript_args)r*  r  r  r  r  r  r5   r  r  r  	sysconfigget_config_varr  shutilwhichr@  r   r  
check_call
setuptools	Extensionr   rq   r.  setup)r)   r  srcdirZcuda_lib_dirsZ	cuda_pathZcu_include_dirr  Ztriton_include_dirZcuda_headerZtriton_cuda_headerr  soccr2  r3  Zpy_include_dircc_cmdrG   r7  r9  r6  r:  r8  extr]   r,   r,   r-   _buildW  sf    



,


0rK  c                 C  sF   dd |  D }|  dd|  | }t|d }|S )Nc                 S  s&   i | ]\}}||d  dkrdn|qS )r   r   ptrr,   r  r,   r,   r-   rO    r<   z%make_so_cache_key.<locals>.<dictcomp>-r  r  )r  r5   r   r  md5r  r  )Zversion_hashr  rA   r  r,   r,   r-   make_so_cache_key  s    rO  c                   sb   dd   fdd|D }|  dd |  d| d| d| d| }t|d }|S )Nc                 S  s   t | jt | jfS rM   rE   r  r  confr,   r,   r-   r    r<   z#make_fn_cache_key.<locals>.<lambda>c                   s   g | ]} |qS r,   r,   r:   rR  Zget_conf_keyr,   r-   r;     r<   z%make_fn_cache_key.<locals>.<listcomp>rM  r  r  )r5   r   r  rN  r  r  )Zfn_hashr  configsrA   r  r  configs_keyr  r,   rT  r-   make_fn_cache_key  s
    0rW  zCallable[[str], bytes]r   )run_if_foundrun_if_not_foundc                 C  s   | dd }|st| |rt|| |}t|tr8|nt|d}t|	 }	|of|	|d | k}
||	|
dfS | }t|tr|nt|d}t|	 }	| 
||t|trdn| ||	ddfS )NrB   r   r  rN  TF)r  r  r  rh   r   r4   r  r  rN  r  r(  )Zcache_managerZforce_compile	file_namemetadatarX  rY  r  rm   r$  rN  Zhas_changedr,   r,   r-   read_or_execute  s    r\  c              	   C  s   t tjj ||}t|}|  d}||st }t	||}t
j|d}t|d}	|	| W d    n1 s|0    Y  t| ||}
t|
d$}	|j|	 |dd W d    n1 s0    Y  W d    n1 s0    Y  ||S )Nz.somain.cr  rbTr%  )rO  r    rT  jitversion_keyr  r  tempfileTemporaryDirectoryr  r  r  r5   r!  r"  rK  r(  readr  )r)   r  rA   Zso_cache_keyZso_cache_managerZso_nametmpdirr  src_pathr  rG  r,   r,   r-   	make_stub  s    



(Prg  c                 C  s*   t d| }|d ur&dt|d S | S )Nz!tt\.ptr<(.*)>r   r   )r  r  convert_type_reprr  )r   matchr,   r,   r-   rh    s    rh  c           	        s   t | tjjr|d }|d }|dt }|dd}|dd}dd	   fd
d|D }| j dd|  d| d| d| d| }t	
|d S t | tsJ t	
t|  tjj  d S )NrU  r  rA   r     r  rn  c                 S  s   t | jt | jfS rM   rP  rQ  r,   r,   r-   r    r<   zmake_hash.<locals>.<lambda>c                   s   g | ]} |qS r,   r,   rS  rT  r,   r-   r;     r<   zmake_hash.<locals>.<listcomp>rM  r  r  )rh   r    rT  rU  r  rq   	cache_keyr5   r   r  rN  r  r  r4   r   	read_textr`  ra  )	r   r^   rU  r  rA   r  r  rV  r  r,   rT  r-   	make_hash  s    2rm  z\^\s*func\s+(?:public\s+)?(@\w+)(\((?:%\w+: \S+(?: \{\S+ = \S+ : \S+\})?(?:, )?)*\))\s*\{\s*$z=\.(?:visible|extern)\s+\.(?:entry|func)\s+(\w+)\s*\(([^)]*)\))ttirttgirr  z-%\w+: ([^,^\)\s]+)(?: \{\S+ = \S+ : \S+\})?,?z\.param\s+\.(\w+)c              	     s  | dd   d u r>tj }tj|  d d  d   tj t }| dt | dd| d d	krzd
nd| dt fddd ffddfddffdd fddfdd  fddfdd  fddfdd  fddfd}t	t
jjr| dd |d d u rDt gtdksVJ |d< j}d}t	trdd tdD |d< nt	tsJ tjd\}}t }	dd l}
|
t| |	|
j}|d|d }|
t| }d d! |D }d"d t|D t|  !|}t"|}t#t$fi |}t	t
jjrzjd# }}ntjd\}}d }|%| d$rt&|'| d$}t()|}W d    n1 s0    Y  n6t d%}|d&krd'|v sJ d(|d' |d'< t|  !|}t }}t|* |d  D ]\}\}}|'| d| }||kr|}n^tj+|r||d) v rtj,||d) | kr||}n||}|-|| d|  tj+|rtj,||d) |< |d*kr|nt|||< |d+kr@d'|vr@t.||d'< |d&krVt/||d,< |}qN|j-t(0|| d$d-d. t1|||S )/NrH  r   r  r   rA   r  rj  r  K   rn  r   r  c                   s    S rM   r,   r  r  r,   r-   r    r<   zcompile.<locals>.<lambda>c                   s
   t |  S rM   r  rq  ry   r,   r-   r    r<   c                   s   t |  d S r  )r  r  )rU  rA   r  r,   r-   r    r<   c                   s
   t |  S rM   rr  rq  rs  r,   r-   r    r<   c                   s   t |  S rM   )r  rt  )
capabilityr  r  r,   r-   r    r<   c                 S  s   t |  S rM   r   rl  rq  r,   r,   r-   r    r<   c                   s   t |  S rM   )r  rt  )ru  r  r,   r-   r    r<   c                 S  s   t |  S rM   rv  rq  r,   r,   r-   r    r<   c                   s
   t |  S rM   )r  rt  ru  r,   r-   r    r<   c                 S  s   t |  S rM   )r   
read_bytesrq  r,   r,   r-   r    r<   c                   s
   t |  S rM   )r  rt  rw  r,   r-   r    r<   )r   rn  ro  llirr  cubinrU  r  c                 S  s   i | ]\}}||  qS r,   r  r  r,   r,   r-   rO  #  r<   zcompile.<locals>.<dictcomp>r  rB   c                 S  s   g | ]}t |qS r,   )rh  r9   r,   r,   r-   r;   /  r<   zcompile.<locals>.<listcomp>c                 S  s   i | ]\}}||qS r,   r,   r  r,   r,   r-   rO  0  r<   r   z.json)r  r  ctimer  sharedz/ptx compilation must provide shared memory sizer{  rz  ry  r)   Fr_  )2r  torchr  current_deviceZget_device_capabilityrj   rk   ry   rq   rh   r    rT  rU  r  r   r`   r4   r   r  r  r  basenamer   rl  r  r  prototype_patternr  r  findallarg_type_patternr   r   r  rg  r  rm  r  r!  r  jsonloadr  r  getctimer(  Zget_shared_memory_sizer  dumpsCompiledKernel)r   r^   deviceasmZstagesr)   Zfirst_stager0   rk   r  r  ri  typesZ	param_tysso_pathZfn_cache_managerrJ  r[  r  rm   r[  compiler  Znext_moduler,   )	ru  rU  rA   ry   r  r   r  r  r  r-   r    s    







,
"




r  c                      sF   e Zd ZdZdZdd Zdd Z fddZdd	 Zdd
dZ	  Z
S )r  Nc                 C  st   dd l }|jd|}|j|}|j| t|d| _|d | _|d | _	|d | _
|| _|| _d | _d | _d S )Nr   Z__triton_launcherZlaunchr|  r  r  )importlib.utilutilspec_from_file_locationmodule_from_specloaderexec_moduleri   	c_wrapperr|  r  r  r  r[  	cu_modulecu_function)rN   r  r[  r  	importlibspecr  r,   r,   r-   rO   k  s    


zCompiledKernel.__init__c                 C  sz   | j d urd S tj }t  t|d }| j|krDt| j|dt	| j
d | jd | j|\}}}}|| _ || _d S )NZmax_shared_memzshared memoryr)   rz  )r  r}  r  r~  init_cuda_utils
cuda_utilsget_device_propertiesr|  r  load_binaryr[  r  r  )rN   r  Z
max_sharedr  r;  Zn_regsZn_spillsr,   r,   r-   _init_handles  s    


&zCompiledKernel._init_handlesc                   s   |dkr|    t |S )Nr  )r  r  __getattribute__)rN   r)   r  r,   r-   r    s    zCompiledKernel.__getattribute__c                   s       d d fdd
}|S )N)streamc                   sR   | d u rt j j} j d  d  d jj| jtj	tj
g
|R   d S )Nr   r   r   )r}  r  Zcurrent_streamZcuda_streamr  r  r|  r  r  launch_enter_hooklaunch_exit_hook)r  r]   gridrN   r,   r-   runner  s    $
z*CompiledKernel.__getitem__.<locals>.runner)r  )rN   r  r  r,   r  r-   r6    s    zCompiledKernel.__getitem__c              	   C  s   d| j v r| j d S t \}}zTt|d }|| j d  W d    n1 sR0    Y  t||| _W t| nt| 0 | j| j d< | jS )Nsassr  rz  )	r  rb  mkstempr!  r"  r   r  r  remove)rN   Zfunfdr  rz  r,   r,   r-   get_sass  s    

.zCompiledKernel.get_sass)N)r`   ra   rb   r  r  rO   r  r  r6  r  r  r,   r,   r  r-   r  e  s   
r  c                      s0   e Zd Z fddZedd Zdd Z  ZS )	CudaUtilsc                   s"   t | dstt| | | _| jS )Ninstance)r]  r  r  __new__r  )clsr  r,   r-   r    s    
zCudaUtils.__new__c                   C  s   dS )Na  
        #include <cuda.h>

        #include "cuda.h"
        #define PY_SSIZE_T_CLEAN
        #include <Python.h>

        static inline void gpuAssert(CUresult code, const char *file, int line)
        {
           if (code != CUDA_SUCCESS)
           {
              const char* prefix = "Triton Error [CUDA]: ";
              const char* str;
              cuGetErrorString(code, &str);
              char err[1024] = {0};
              strcat(err, prefix);
              strcat(err, str);
              PyErr_SetString(PyExc_RuntimeError, err);
           }
        }

        #define CUDA_CHECK(ans) { gpuAssert((ans), __FILE__, __LINE__); if(PyErr_Occurred()) return NULL; }

        static PyObject* getDeviceProperties(PyObject* self, PyObject* args){
            int device_id;
            if(!PyArg_ParseTuple(args, "i", &device_id))
                return NULL;
            // Get device handle
            CUdevice device;
            cuDeviceGet(&device, device_id);

            // create a struct to hold device properties
            int max_shared_mem;
            int multiprocessor_count;
            int sm_clock_rate;
            int mem_clock_rate;
            int mem_bus_width;
            CUDA_CHECK(cuDeviceGetAttribute(&max_shared_mem, CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN, device));
            CUDA_CHECK(cuDeviceGetAttribute(&multiprocessor_count, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, device));
            CUDA_CHECK(cuDeviceGetAttribute(&sm_clock_rate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, device));
            CUDA_CHECK(cuDeviceGetAttribute(&mem_clock_rate, CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE, device));
            CUDA_CHECK(cuDeviceGetAttribute(&mem_bus_width, CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH, device));


            return Py_BuildValue("{s:i, s:i, s:i, s:i, s:i}", "max_shared_mem", max_shared_mem,
                                       "multiprocessor_count", multiprocessor_count,
                                       "sm_clock_rate", sm_clock_rate,
                                       "mem_clock_rate", mem_clock_rate,
                                       "mem_bus_width", mem_bus_width);
        }

        static PyObject* loadBinary(PyObject* self, PyObject* args) {
            const char* name;
            const char* data;
            Py_ssize_t data_size;
            int shared;
            int device;
            if(!PyArg_ParseTuple(args, "ss#ii", &name, &data, &data_size, &shared, &device)) {
                return NULL;
            }
            CUfunction fun;
            CUmodule mod;
            int32_t n_regs = 0;
            int32_t n_spills = 0;
            // create driver handles
            CUDA_CHECK(cuModuleLoadData(&mod, data));
            CUDA_CHECK(cuModuleGetFunction(&fun, mod, name));
            // get allocated registers and spilled registers from the function
            CUDA_CHECK(cuFuncGetAttribute(&n_regs, CU_FUNC_ATTRIBUTE_NUM_REGS, fun));
            CUDA_CHECK(cuFuncGetAttribute(&n_spills, CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, fun));
            n_spills /= 4;
            // set dynamic shared memory if necessary
            int shared_optin;
            CUDA_CHECK(cuDeviceGetAttribute(&shared_optin, CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN, device));
            if (shared > 49152 && shared_optin > 49152) {
              CUDA_CHECK(cuFuncSetCacheConfig(fun, CU_FUNC_CACHE_PREFER_SHARED));
              int shared_total, shared_static;
              CUDA_CHECK(cuDeviceGetAttribute(&shared_total, CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR, device));
              CUDA_CHECK(cuFuncGetAttribute(&shared_static, CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, fun));
              CUDA_CHECK(cuFuncSetAttribute(fun, CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, shared_optin - shared_static));
            }

            if(PyErr_Occurred()) {
              return NULL;
            }
            return Py_BuildValue("(KKii)", (uint64_t)mod, (uint64_t)fun, n_regs, n_spills);
        }

        static PyMethodDef ModuleMethods[] = {
          {"load_binary", loadBinary, METH_VARARGS, "Load provided cubin into CUDA driver"},
          {"get_device_properties", getDeviceProperties, METH_VARARGS, "Get the properties for a given device"},
          {NULL, NULL, 0, NULL} // sentinel
        };

        static struct PyModuleDef ModuleDef = {
          PyModuleDef_HEAD_INIT,
          "cuda_utils",
          NULL, //documentation
          -1, //size
          ModuleMethods
        };

        PyMODINIT_FUNC PyInit_cuda_utils(void) {
          PyObject *m = PyModule_Create(&ModuleDef);
          if(m == NULL) {
            return NULL;
          }
          PyModule_AddFunctions(m, ModuleMethods);
          return m;
        }
        r,   r,   r,   r,   r-   _generate_src  s    zCudaUtils._generate_srcc              	   C  s.  |   }t|d }t|}d}||st }t	j
|d}t|d}|| W d    n1 st0    Y  td||}t|d$}|j| |dd W d    n1 s0    Y  W d    n1 s0    Y  d	d l}	|	jd||}
|	j|
}|
j| |j| _|j| _d S )
Nr  zcuda_utils.sor]  r  r  r^  Tr_  r   )r  r  rN  r  r  r  r  rb  rc  r  r  r5   r!  r"  rK  r(  rd  r  r  r  r  r  r  r  r  r  )rN   r  r  r  fnamere  rf  r  rG  r  r  r  r,   r,   r-   rO   #  s$    

(PzCudaUtils.__init__)r`   ra   rb   r  staticmethodr  rO   r  r,   r,   r  r-   r    s   
qr  c                   C  s   t d u rt a d S rM   )r  r  r,   r,   r,   r-   r  8  s    r  )N)NN)^
__future__r   r   
contextlib	functoolsr  r,  r  r  r  r@  r  rX  r>  rb  rz  collectionsr   pathlibr   r   typingr   r   r   r	   r
   rC  r}  Zfilelockr   r    Ztriton._C.libtriton.tritonZ_CZ	libtritonrj   r  r   Ztools.disasmr   r   r3   rH   rI   r   rJ   r  r  r  r  r  r  r  r  r  r  r  r  r  r  	lru_cacher  r  setr  r  r  r  r  r  r  r  r*  contextmanagerr.  rK  rO  rW  r\  rg  rh  rm  Zmlir_prototype_patternZptx_prototype_patternr  Zmlir_arg_type_patternZptx_arg_type_patternr  r  r  objectr  r  r  r,   r,   r,   r-   <module>   s   
     v" 7'

	>
  jE 