a
    zd                  	   @  s  d dl mZ d dlmZmZmZ ddlmZ d dl	m
Z
 G dd deZdd	d
dddZdd	d
dddZddddddZdddddddZdddddddZdd
d
d	ddd d!Zd
d
d	d
d"d#d$Zd
d
d	d
d"d%d&Zd
d
d	d
d"d'd(Zd
d
d	d
d"d)d*Zd
d
d	d
d"d+d,Zd
d
dd	d
d-d.d/Zd
d
d	d
d"d0d1Zd
d
d	dd"d2d3Zd
d
d	d
d"d4d5Zd
d
d	d
d"d6d7Zd
d
d	d
d"d8d9Zd
d
d	d
d"d:d;Zd
d
d	d
d"d<d=Zd
d	d>d?d@Z d
d
d	d
d"dAdBZ!d
d
d	d
d"dCdDZ"d
d
d	d
d"dEdFZ#d
d
dGdHdIZ$d
d	d
dJdKdLZ%d
d
d
dJdMdNZ&d
dOdPdQdRZ'd
d
d	d
d"dSdTZ(d
d
d	d
d"dUdVZ)d
d
d	d
d"dWdXZ*d
d
d	d
d"dYdZZ+d
d
d	d
d"d[d\Z,d
d
d	d
d"d]d^Z-ddd	d
d_d`daZ.dbdd	d
dcdddeZ/d
dbd	d
dfdgdhZ0d
dd	d
didjdkZ1d
d
dd	d
dldmdnZ2d
d	d
dJdodpZ3d
dbd	d
dqdrdsZ4d
d
d	d
ddtduZ5d
dd	d
dvdwdxZ6d
dd	d
dvdydzZ7d
d{d{d|d|dd	d
d}d~dZ8d
d
d{d	d
dddZ9d
d
d
d	d
dddZ:d
d
d
d|d	ddddZ;d
d
d
d	d
dddZ<d
d
d
d	d
dddZ=d
d
d
d	d
dddZ>d
d
d
d	d
dddZ?d
d
d
d	d
dddZ@d
d
d
d	d
dddZAd
d
d
d	d
dddZBd
d
dd	d
dddZCd
d
d
d	d
dddZDd
dd	d|ddd
dddZEd
dd	d
diddZFd
dd	d
diddZGd
dd	d
diddZHd
dd	d
diddZId
dd	d
diddZJd
dd	d
diddZKd
d
d	d
dddZLd
d	d
dddZMd
d	d
dddZNd
d	d
dddZOd
d	d
dddZPd
d	d
dddZQd
d	d
dddZRd
dbd
dddZSd
dbd
dddZTd	d
dÜddńZUd|dd	d
dǜddɄZVdS )    )annotations)ListOptionalTuple   )core)irc                      s   e Zd Z fddZ  ZS )IncompatibleTypeErrorImplc                   s@   || _ || _d| j   d | j  | _tt| | j d S )Nzinvalid operands of type  and )type_atype_b__repr__messagesuperr	   __init__)selfr   r   	__class__ a/var/www/html/stable-diffusion-webui/venv/lib/python3.9/site-packages/triton/language/semantic.pyr      s    z"IncompatibleTypeErrorImpl.__init__)__name__
__module____qualname__r   __classcell__r   r   r   r   r	   
   s   r	   intz
ir.builderz	tl.tensor)axisbuilderreturnc                 C  s   t || t jS N)tltensorZcreate_get_program_idint32r   r   r   r   r   
program_id   s    r#   c                 C  s   t || t jS r   )r   r    Zcreate_get_num_programsr!   r"   r   r   r   num_programs   s    r$   ztl.dtype)a_tyb_tyr   c                 C  sx   | j }|j }| j}|j}||kr0||kr,| S |S |tjjjkrN||krJ| S |S |tjjjkrl||krh|S | S dstJ d S NF)int_bitwidthint_signednessr   dtypeZ
SIGNEDNESSZUNSIGNED)r%   r&   Za_rankZb_rankZa_snZb_snr   r   r   integer_promote_impl"   s    r+   bool)r%   r&   
div_or_modr   c                 C  s   |   s|  rtjS |  s&| r,tjS |  s<| rL|rFtjS tjS |  s\| r|rftjS |  r|| r|tjS tjS | 	 r|	 sdsJ |r| j
|j
krtd|   d |  d t| |S )NFzCannot use /, #, or % with r
   x because they have different signedness;this is unlikely to result in a useful answer. Cast them to the same signedness.)Zis_fp64r   float64is_fp32float32is_fp16float16is_bf16bfloat16is_intr)   
ValueErrorr   r+   )r%   r&   r-   r   r   r   computation_type_impl2   s&     r8   None)r   r   allow_ptr_ar   c                 C  sF   |   rB|st| ||  r0| |kr0t| || rBt| |d S r   )is_ptrr	   is_floating)r   r   r:   r   r   r   check_ptr_type_implX   s    

r=   FTzTuple[tl.tensor, tl.tensor])lhsrhsr   r   c           
      C  sx   t | ||\} }| jj}|jj}t||| t||| |rp| sp| spt|||}	t| |	|} t||	|}| |fS r   )broadcast_impl_valuetypescalarr=   r;   r8   cast)
r>   r?   r   Zallow_lhs_ptrZallow_rhs_ptrZarithmetic_checkr-   Z
lhs_sca_tyZ
rhs_sca_ty
ret_sca_tyr   r   r   binary_op_type_checking_impld   s    rE   )inputotherr   r   c                 C  s   t | ||dd\} }| jj}|jj}| r>| s>||  } }| r`t|| j|j| jS | rt|	| j|j| jS |
 rt|| j|j| jS dsJ d S NTF)rE   rA   rB   r;   r   r    create_addptrhandler<   Zcreate_faddr6   Z
create_addrF   rG   r   input_scalar_tyother_scalar_tyr   r   r   addx   s    
rN   c                 C  s   t | ||dd\} }| jj}| rDt|| jt||j| jS |	 rft|
| j|j| jS | rt|| j|j| jS dsJ d S rH   )rE   rA   rB   r;   r   r    rI   rJ   minusr<   Zcreate_fsubr6   Z
create_subrF   rG   r   	scalar_tyr   r   r   sub   s    rR   c                 C  sh   t | ||\} }| jj}| r:t|| j|j| jS | r\t|	| j|j| jS dsdJ d S r'   )
rE   rA   rB   r<   r   r    Zcreate_fmulrJ   r6   Z
create_mulrP   r   r   r   mul   s    rS   c                 C  s   t | ||dddd\} }| jj}|jj}| rF| rFt|||}n| rd| rdt| ||} nn| r| rt| tj|} t|tj|}n@| r| r|j|jkrt|||}qt| ||} ndsJ t	|
| j|j| jS NFT)rE   rA   rB   r<   r6   rC   r   r1   Zfp_mantissa_widthr    create_fdivrJ   rK   r   r   r   truediv   s     rV   c                 C  s   t | ||dddd\} }| jj}|jj}| r| rt||}t| ||} t|||}| r|t|	| j
|j
| jS t|| j
|j
| jS dsJ d S rT   )rE   rA   rB   r6   r+   rC   is_int_signedr   r    Zcreate_sdivrJ   Zcreate_udiv)rF   rG   r   rL   rM   ret_tyr   r   r   floordiv   s    
rY   )rF   rG   ieee_roundingr   r   c                 C  s^   | j j}|j j}| r | s(tdt| ||dddd\} }|| j|j}t|| j S )Nz4both operands of fdiv must have floating scalar typeFT)	rA   rB   r<   r7   rE   rU   rJ   r   r    )rF   rG   rZ   r   rL   rM   retr   r   r   fdiv   s    r\   c              	   C  s   t | ||dddd\} }| jj}|jj}| rXt| ttt| |d|||||}|S | r|j	|j	krt
d|  d |  d | rt|| j|j| jS t|| j|j| jS dsJ d S )NFTzCannot mod z by r.   )rE   rA   rB   r<   rR   rS   floorr\   r6   r)   r7   r   rW   r   r    Zcreate_sremrJ   Zcreate_urem)rF   rG   r   rQ   rM   r[   r   r   r   mod   s"     r^   c                 C  sz   t | ||ddd\} }| jj}|jj}| r6| s@t||t||}||kr^t| ||} ||krrt|||}| |fS r'   )rE   rA   rB   r6   r	   r+   rC   )rF   rG   r   input_sca_tyZother_sca_tyrD   r   r   r   bitwise_op_type_checking_impl  s    

r`   c                 C  s*   t | ||\} }t|| j|j| jS r   )r`   r   r    Z
create_andrJ   rA   rF   rG   r   r   r   r   and_  s    rb   c                 C  s*   t | ||\} }t|| j|j| jS r   )r`   r   r    Z	create_orrJ   rA   ra   r   r   r   or_  s    rc   c                 C  s*   t | ||\} }t|| j|j| jS r   )r`   r   r    Z
create_xorrJ   rA   ra   r   r   r   xor_$  s    rd   c                 C  sD   | j  st| td|} |j  s8t|td|}t| ||S Nint1)rA   is_int1bitcastr   r*   rb   ra   r   r   r   logical_and+  s
    

ri   c                 C  sD   | j  st| td|} |j  s8t|td|}t| ||S re   )rA   rg   rh   r   r*   rc   ra   r   r   r   
logical_or3  s
    

rj   rF   r   c                 C  s&   | j  st| td|} t| |S re   )rA   rg   rh   r   r*   invertrk   r   r   r   not_;  s    
rm   c                 C  s*   t | ||\} }t|| j|j| jS r   )r`   r   r    Zcreate_lshrrJ   rA   ra   r   r   r   lshrA  s    rn   c                 C  s*   t | ||\} }t|| j|j| jS r   )r`   r   r    Zcreate_ashrrJ   rA   ra   r   r   r   ashrH  s    ro   c                 C  s*   t | ||\} }t|| j|j| jS r   )r`   r   r    Z
create_shlrJ   rA   ra   r   r   r   shlO  s    rp   )rF   r   c                 C  s   | S r   r   )rF   r   r   r   plusZ  s    rq   )rF   r   r   c                 C  sH   | j j}| r$td|  d t||||}t	|| |S )Nz$wrong type argument to unary minus ())
rA   rB   r;   r7   r   r   r    get_null_valueto_irrR   )rF   r   r_   _0r   r   r   rO   ^  s
    rO   c                 C  sP   | j j}| s| r,td|  d t||	||}t
| ||S )Nz%wrong type argument to unary invert (rr   )rA   rB   r;   r<   r7   r   r   r    Zget_all_ones_valuert   rd   )rF   r   r_   Z_1r   r   r   rl   g  s
    rl   ztl.block_type)vr   c                 C  s&   | j  stjS | j j}ttj|S r   )rA   is_blockr   rf   shape
block_type)rv   rx   r   r   r   
_bool_likes  s    
rz   c                 C  s   t | ||\} }| jj}| r<t|| j|jt| S |	 r|
 rht|| j|jt| S t|| j|jt| S dsJ d S r'   )rE   rA   rB   r<   r   r    Zcreate_fcmpOGTrJ   rz   r6   rW   Zcreate_icmpSGTZcreate_icmpUGTrP   r   r   r   greater_thanz  s    r{   c                 C  s   t | ||\} }| jj}| r<t|| j|jt| S |	 r|
 rht|| j|jt| S t|| j|jt| S dsJ d S r'   )rE   rA   rB   r<   r   r    Zcreate_fcmpOGErJ   rz   r6   rW   Zcreate_icmpSGEZcreate_icmpUGErP   r   r   r   greater_equal  s    r|   c                 C  s   t | ||\} }| jj}| r<t|| j|jt| S |	 r|
 rht|| j|jt| S t|| j|jt| S dsJ d S r'   )rE   rA   rB   r<   r   r    Zcreate_fcmpOLTrJ   rz   r6   rW   Zcreate_icmpSLTZcreate_icmpULTrP   r   r   r   	less_than  s    r}   c                 C  s   t | ||\} }| jj}| r<t|| j|jt| S |	 r|
 rht|| j|jt| S t|| j|jt| S dsJ d S r'   )rE   rA   rB   r<   r   r    Zcreate_fcmpOLErJ   rz   r6   rW   Zcreate_icmpSLEZcreate_icmpULErP   r   r   r   
less_equal  s    r~   c                 C  sl   t | ||\} }| jj}| r<t|| j|jt| S |	 r`t|
| j|jt| S dshJ d S r'   )rE   rA   rB   r<   r   r    Zcreate_fcmpOEQrJ   rz   r6   Zcreate_icmpEQrP   r   r   r   equal  s    r   c                 C  sl   t | ||\} }| jj}| r<t|| j|jt| S |	 r`t|
| j|jt| S dshJ d S r'   )rE   rA   rB   r<   r   r    Zcreate_fcmpUNErJ   rz   r6   Zcreate_icmpNErP   r   r   r   	not_equal  s    r   )startendr   r   c                 C  sH   t | trt |tstd||  g}ttj|}t|| ||S )Nz/arange's arguments must be of type tl.constexpr)
isinstancer   r7   r   ry   r!   r    Zcreate_make_range)r   r   r   rx   rX   r   r   r   arange  s
    
r   z	List[int])rx   r*   r   r   c                 C  sT   |dkr| ||}nt|d|j }||}t|| }t||| |S )Nr   get_)rs   rt   getattrnamer   ry   r    create_splat)rx   valuer*   r   _valueZget_value_fnrX   r   r   r   full  s    r   )rF   	dst_shaper   r   c                 C  sP   d}|D ]}||9 }q| j j|kr*tdt| j j|}t|| j||S )Nr   z$cannot view block of different shape)	rA   numelr7   r   ry   rB   r    Zcreate_viewrJ   )rF   r   r   r   srX   r   r   r   view  s    
r   )rF   r   r   r   c                 C  s>   t | jj}||d t| jj|}t|| j	||S )Nr   )
listrA   rx   insertr   ry   rB   r    create_expand_dimsrJ   )rF   r   r   r   rX   r   r   r   expand_dims  s    r   )r>   r?   can_reorderr   r   c                 C  sX   |sJ dt | jdksJ t| jj| jd |jd  g}t|| j|j|S )Nz;current implementation of `cat` always may reorder elementsr   r   )	lenrx   r   ry   rA   rB   r    Z
create_catrJ   )r>   r?   r   r   ret_typer   r   r   cat  s    "r   c                 C  sJ   t | jdkrtdt| jj| jd | jd g}t|| j	|S )N   z!Only 2D tensors can be transposedr   r   )
r   rx   r7   r   ry   rA   rB   r    Zcreate_transrJ   )rF   r   r   r   r   r   trans  s     r   )rF   rx   r   r   c                 C  s   | j  s.t| j |}t|| j||S | j  }t|t|kr\t	d| d| ||krh| S t
|D ]F\}}|| |krp|dkrpt	d||  d| d| d| d| 
qpt| j j|}t|| j||S )Nz!Cannot broadcast, rank mismatch: z, r   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension : )rA   rw   r   ry   r    r   rJ   get_block_shapesr   r7   	enumeraterB   create_broadcast)rF   rx   r   rX   Z	src_shapeiitemr   r   r   broadcast_impl_shape  s*    


r   c              	   C  sp  | j }|j }| rJ| sJt|j|j}t||j|	 |}n| s| rt|j|j}t|| j|	 |} n| rh| rh|	 }|	 }t
|t
|k rtt
|t
|D ]8}t|| j|t|j|dg } | j }|	 }qn`t
|t
|k rltt
|t
|D ]:}t||j|t|j|dg }|j }|	 }q0t
|t
|ksJ g }t|D ]|\}	}
||	 }|
dkr|| nT|dkr||
 n>|
|kr||
 n(tdt|	 d t|
 d t| q||kr:t|j|}t|| j||} ||krht|j|}t||j||}| |fS )Nr   z?Cannot make_shape_compatible: incompatible dimensions at index r   r
   )rA   rw   r   ry   rB   rx   r    r   rJ   r   r   ranger   r   appendr7   strr   )r>   r?   r   Zlhs_tyZrhs_tyZ	lhs_shapeZ	rhs_shapeZdim	ret_shaper   leftrightrX   r   r   r   r@   .  sb    &&





r@   )rF   dst_tyr   r   c                 C  s   | j }| r t|| j  }||kr,| S |j}|j}| sH| rTt| ||S |j}|j}||krt	dt
| d t
| t|| j|||S )Nz!Cannot bitcast data-type of size zto data-type of size )rA   rw   r   ry   r   rB   r;   rC   primitive_bitwidthr7   r   r    create_bitcastrJ   rt   )rF   r   r   src_ty
src_sca_ty
dst_sca_tyZsrc_bitsZdst_bitsr   r   r   rh   i  s$    rh   c                 C  s  | j }| r t|| j  }||kr,| S |j}|j}| rH| sX| rt| rtt|	| j
|||S | r| r| r| stt| tj|||S | o| o|j|jk}|rt|| j
|||S | o| o|j|jk }|r(t|| j
|||S | r| r|j|jksX|j|jkr| oj|  }| r| j|}	t||	| j}
t| |
|S t|| j
||||S | rL| rL| r
| j|}	t||	| j}
t| |
|S | r0t|| j
|||S t|| j
|||S | r| r| st| st| | j
|||S t|!| j
|||S |" r| r|j}|dkrt|#| j
|||S |dkrtt| tj$|t|%dtj$|S | rL|" rLt|&| j
|||S |" r||" r|t|'| j
|||S dsJ d|  d| d S )N@   r   r   Fzcannot cast z to )(rA   rw   r   ry   r   rB   Zis_customized_floatingr<   r    Zcreate_fp_to_fprJ   rt   r2   r0   r4   rC   r1   r   Zcreate_fp_truncZcreate_fp_extr6   r(   r)   rW   is_boolr*   rs   r   Zcreate_int_castZis_standard_floatingZcreate_fp_to_siZcreate_fp_to_uiZcreate_ui_to_fpZcreate_si_to_fpr;   Zcreate_ptr_to_intint64Z	get_int64Zcreate_int_to_ptrr   )rF   r   r   r   r   r   Ztruncate_fpZext_fpZsign_extendtyru   Zbitwidthr   r   r   rC     s    










rC   zOptional[tl.tensor]r   )ptrmaskrG   cache_modifiereviction_policyis_volatiler   r   c              
   C  s  | j j std| j   | j  rT|r>t|| j  |}|rTt|| j  |}| j j}|j}|t	j
krt	j}t	||j}t| ||} |rt|||}tjj}	|r|dkrtjj}	n"|dkrtjj}	ntd| dtjj}
|r |dkrtjj}
n$|dkrtjj}
ntd| d| j  rD| j  }t	||}n|}|sv|r\td	t	|| j|	|
||S t	|| j|j|r|jnd |	|
||S d S )
Nz(Pointer argument of load instruction is z.caz.cgzCache modifier z not supportedZ
evict_lastZevict_firstzEviction policy z)`other` cannot be provided without `mask`)rA   rB   r;   r7   r   rw   r   r   
element_tyr   rf   int8pointer_typeaddress_spacerC   r   ZCACHE_MODIFIERNONECAZCGZEVICTION_POLICYNORMALZ
EVICT_LASTZEVICT_FIRSTry   r    Zcreate_loadrJ   Zcreate_masked_load)r   r   rG   r   r   r   r   ptr_tyelt_tycacheZevictionrx   r   r   r   r   load  s\    







r   )r   valr   r   r   c                 C  s   | j j std| j   | j  r:t|| j  |}|rZ| j  rZt|| j  |}| j j}|j}|t	j
krt	j}t	||j}t| ||} t|||}|st	|| j|jt	jS |j j stdt	|| j|j|jt	jS )N)Pointer argument of store instruction is z"Mask must have boolean scalar type)rA   rB   r;   r7   r   rw   r   r   r   r   rf   r   r   r   rC   r    Zcreate_storerJ   voidr   Zcreate_masked_store)r   r   r   r   r   r   r   r   r   store0  s$    

r   )r   cmpr   r   r   c                 C  s:   | j jj}|jdvrtdt|| j|j|j|j S )N)       r   z9atomic_cas only supports elements with width {16, 32, 64})	rA   rB   r   r   r7   r   r    Zcreate_atomic_casrJ   )r   r   r   r   r   r   r   r   
atomic_casO  s    

r   z&Tuple[tl.tensor, tl.tensor, tl.tensor])r   r   r   opr   r   c                 C  s   | j j std| j   | j jj}|tju rJ|dkrJtd| d |tjtj	tj
tjfv rztd| d t| | j  r|rt|| j  |}|rt|| j  |}t|| j jj|}|s|d}tj}| j  r
||| j  }ttj| j  }t||}| ||fS )Nr   rN   Zatomic_z does not support fp16z does not support T)rA   rB   r;   r7   r   r   r   r3   rf   r   int16r5   r   rw   r   r   rC   Zget_int1r   ry   r    )r   r   r   r   r   r   Zmask_irZmask_tyr   r   r   atom_red_typechecking_implY  s*    


r   c                 C  s4  t | ||d|\} }}|jj}| rv| rRt|tj	j
| j|j|j|jS t|tj	j| j|j|j|jS t|tj|}t| ttjd|}t|t|d||}t|t|d||}t|tj	j
|j|jt|||j|j}	t|tj	j|j|jt|||j|j}
t||	|
|S )Nmaxr   r   )r   rA   rB   r6   rW   r   r    create_atomic_rmwr   	ATOMIC_OPMAXrJ   UMAXrh   r!   r   r|   get_fp32r}   rb   UMINwherer   r   r   r   sca_tyZi_valZi_ptrposnegZpos_retZneg_retr   r   r   
atomic_maxv  s2    ,,r   c                 C  s4  t | ||d|\} }}|jj}| rv| rRt|tj	j
| j|j|j|jS t|tj	j| j|j|j|jS t|tj|}t| ttjd|}t|t|d||}t|t|d||}t|tj	j
|j|jt|||j|j}	t|tj	j|j|jt|||j|j}
t||	|
|S )Nminr   r   )r   rA   rB   r6   rW   r   r    r   r   r   MINrJ   r   rh   r!   r   r|   r   r}   rb   r   r   r   r   r   r   
atomic_min  sJ    r   c                 C  sV   t | ||d|\} }}|jj}| r.tjjntjj}t	|
|| j|j|j|jS )NrN   )r   rA   rB   r<   r   r   FADDADDr   r    r   rJ   )r   r   r   r   r   r   r   r   r   
atomic_add  s    r   c                 C  s:   t | ||d|\} }}t|tjj| j|j|j|jS )Nand)	r   r   r    r   r   r   ANDrJ   rA   r   r   r   r   r   r   r   
atomic_and  s    r   c                 C  s:   t | ||d|\} }}t|tjj| j|j|j|jS )Nor)	r   r   r    r   r   r   ORrJ   rA   r   r   r   r   	atomic_or  s    r   c                 C  s:   t | ||d|\} }}t|tjj| j|j|j|jS )Nxor)	r   r   r    r   r   r   XORrJ   rA   r   r   r   r   
atomic_xor  s    r   c                 C  s:   t | ||d|\} }}t|tjj| j|j|j|jS )NZxchg)	r   r   r    r   r   r   ZXCHGrJ   rA   r   r   r   r   atomic_xchg  s    r   )r>   r?   
allow_tf32r   r   c           	      C  s"  | j  r|j  sJ | j|jks,J dt| jdkrHt|jdksLJ | jd j|jd jkshJ | jd jdkr| jd jdkr|jd jdksJ d| j j r|d}t	j
}n|d}t	j}| j jd }|j jd }||||g}t	|||g}t	|| j|j|||S )Nz%lhs and rhs must have the same dtype!r   r   r   r   zsmall blocks not supported!)rA   rw   r*   r   rx   r   rB   r6   Z	get_int32r   r!   r   r1   r   ry   r    Z
create_dotrJ   )	r>   r?   r   r   ru   Zret_scalar_tyMNrX   r   r   r   dot  s*      

r   )	conditionxyr   r   c                 C  s   t | tj|} | j rHt| ||\} }t|||\}}t| ||\} }t|||dd\}}| j svt| ||\} }|j}t|| j	|j	|j	|S )NT)
rC   r   rf   rA   rw   r@   rE   r    Zcreate_selectrJ   )r   r   r   r   _rX   r   r   r   r   	  s    

r   zir.REDUCE_OP)rF   r   r   r   FLOAT_OPINT_OPr   c                 C  st  | j j}|}| r2|jdkr2t| tj|} tj}|tju rPt| tj|} tj}|	 rt
jjt
jjt
jjt
jjt
jjt
jjt
jjt
jji}||v r|| }|t
jju s|t
jju rtj}n|t
jju s|t
jju rtj}| j j}	g }
t|	D ]\}}||kr|
| q|
rt||
}n|}| rDt|| j|||S | rft|| j|||S dspJ d S )Nr   F)rA   rB   r6   r(   rC   r   r!   r5   r1   Zis_int_unsignedr   	REDUCE_OPr   r   r   r   ARGMINZARGUMINARGMAXZARGUMAXARGFMAXARGFMINrx   r   r   ry   r<   r    Zcreate_reducerJ   )rF   r   r   r   r   r   rQ   Zout_scalar_tyZint_op_to_unitrx   r   r   r   Zres_tyr   r   r   reduce_impl  sB    


r   c                 C  s   t | ||dtjjtjjS )Nr   )r   r   r   ZFMINr   rF   r   r   r   r   r   r   R  s    r   c                 C  s   t | ||dtjjtjjS )Nargmin)r   r   r   r   r   r   r   r   r   r   V  s    r   c                 C  s   t | ||dtjjtjjS )Nr   )r   r   r   ZFMAXr   r   r   r   r   r   Z  s    r   c                 C  s   t | ||dtjjtjjS )Nargmax)r   r   r   r   r   r   r   r   r   r   ^  s    r   c                 C  s   t | ||dtjjtjjS )Nsum)r   r   r   r   r   r   r   r   r   r   b  s    r   c                 C  s2   | j j}| stdt| ||dtjjtjjS )Nz#xor_sum only supported for integersr   )rA   rB   r6   r7   r   r   r   r   )rF   r   r   rQ   r   r   r   xor_sumf  s    r   )r   r   r   r   c                 C  s,   t | ||\} }ddlm} |j| ||dS Nr   )	libdevice)Z_builder)rE    r   Zmulhi)r   r   r   r   r   r   r   umulhiq  s    r  )r   r   r   c                 C  s   ddl m} |j| |dS r   )r  r   r]   )r   r   r   r   r   r   r]   x  s    r]   c                 C  s   t || j| jS r   )r   r    Z
create_exprJ   rA   r   r   r   r   r   exp~  s    r  c                 C  s   t || j| jS r   )r   r    Z
create_logrJ   rA   r  r   r   r   log  s    r  c                 C  s   t || j| jS r   )r   r    Z
create_cosrJ   rA   r  r   r   r   cos  s    r  c                 C  s   t || j| jS r   )r   r    Z
create_sinrJ   rA   r  r   r   r   sin  s    r  c                 C  s   t || j| jS r   )r   r    Zcreate_sqrtrJ   rA   r  r   r   r   sqrt  s    r  )r   valuesr   c                 C  s:   t | jt |krtd| jdt|| j  | S )NzAShape of input to multiple_of does not match the length of valuesztt.divisibilityr   rx   r7   rJ   Zset_attrr   Z	make_attrZget_contextr   r	  r   r   r   multiple_of  s    r  c                 C  s:   t | jt |krtd| jdt|| j  | S )NzDShape of input to max_contiguous does not match the length of valuesztt.contiguityr
  r  r   r   r   max_contiguous  s    r  )r   r   c                 C  s   t |  t jS r   )r   r    Zcreate_barrierr   )r   r   r   r   debug_barrier  s    r  zList[tl.tensor])prefixargsr   r   c                 C  s0   g }|D ]}| |j qt|| |tjS r   )r   rJ   r   r    Zcreate_printfr   )r  r  r   new_argsargr   r   r   printf  s    r  N)FFTF)W
__future__r   typingr   r   r   r  r   r   Ztriton._C.libtriton.tritonr   	Exceptionr	   r#   r$   r+   r8   r=   rE   rN   rR   rS   rV   rY   r\   r^   r`   rb   rc   rd   ri   rj   rm   rn   ro   rp   rq   rO   rl   rz   r{   r|   r}   r~   r   r   r   r   r   r   r   r   r   r@   rh   rC   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  r  r  r  r  r  r   r   r   r   <module>   s   &  		;oB
 (
4