U
    yÊ¹j?þ  ã                
   @  s*  d dl mZ d dlZd dlmZ d dlmZmZmZm	Z	m
Z
 ddlmZ ddlmZ e
d	ƒ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dõ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/d0„Zdddddd1œ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&œd>d?„Z$ddddd&œd@dA„Z%dddBœdCdD„Z&ddddd&œdEdF„Z'ddddd&œdGdH„Z(ddddd&œdIdJ„Z)dddKœdLdM„Z*ddddNœdOdP„Z+ddddNœdQdR„Z,ddSdTœdUdV„Z-ddddd&œdWdX„Z.ddddd&œdYdZ„Z/ddddd&œd[d\„Z0ddddd&œd]d^„Z1ddddd&œd_d`„Z2ddddd&œdadb„Z3dddddcœddde„Z4dfddddgœdhdi„Z5ddfdddjœdkdl„Z6ddfdddjœdmdn„Z7dddddoœdpdq„Z8ddddddrœdsdt„Z9ddddNœdudv„Z:ddfdddwœdxdy„Z;ddddd#œdzd{„Z<ddddd|œd}d~„Z=dd€„ Z>ddddd|œd�d‚„Z?dƒd„„ Z@d…d†„ ZAd‡dˆ„ ZBd‰dŠ„ ZCd‹dŒ„ ZDd�dŽ„ ZEd�d�„ ZFd‘d’„ ZGdd“d“d”d”d”dddd•œ	d–d—„ZHd˜d™„ ZIdšd›„ ZJddd“d”d”dddœœd�dž„ZKdddd”dddŸœd d¡„ZLdddd”dd¢d£œd¤d¥„ZMdddd”ddd¦œd§d¨„ZNdddd”ddd¦œd©dª„ZOdddd”ddd¦œd«d¬„ZPdddd”ddd¦œd­d®„ZQdddd”ddd¦œd¯d°„ZRdddd”ddd¦œd±d²„ZSdddd”ddd¦œd³d´„ZTdddddddµœd¶d·„ZUdddddd¸œd¹dº„ZVd»ddd¼d½œd¾d¿„ZWd»ddd¼d½œdÀdÁ„ZXdÂd	dÃœdÄdÅ„ZYdddddÆœdÇdÈ„ZZeYdÉdÊgdË�ddddÌœdÍdÎ„ƒZ[eYdÉdÊgdË�ddddÌœdÏdÐ„ƒZ\eYdÉdÊgdË�ddddÌœdÑdÒ„ƒZ]eYdÉdÊgdË�ddddÌœdÓdÔ„ƒZ^eYdÉdÊgdË�ddddÌœdÕdÖ„ƒZ_eYdÉdÊgdË�ddddÌœd×dØ„ƒZ`ddddÌœdÙdÚ„ZaddfddÛœdÜdÝ„ZbddfddÛœdÞdß„ZcddfddÛœdàdá„Zddddâœdãdä„Zed”dådddæœdçdè„Zfdd”d”ddddéœdêdë„Zgdìdí„ Zhdödîdï„Ziddddðœdñdò„Zjddddðœdódô„ZkdS )÷é    )ÚannotationsN©Úwraps)ÚListÚOptionalÚSequenceÚTupleÚTypeVaré   )Úiré   )ÚcoreÚTc                      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__© úW/var/www/html/bwave/python/venv/lib/python3.8/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stt‚d S ©NF)Úint_bitwidthÚint_signednessr%   ÚdtypeZ
SIGNEDNESSZUNSIGNEDÚAssertionError)r+   r,   Za_rankZb_rankZa_snZb_snr   r   r   Úinteger_promote_impl'   s    r2   Ú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’| 	¡ sšdsšt
‚|rÊ| j|jkrÊt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%   Zfloat64Úis_fp32Úfloat32Úis_fp16Úfloat16Úis_bf16Úbfloat16Úis_intr1   r/   Ú
ValueErrorr   r2   )r+   r,   r4   r   r   r   Úcomputation_type_impl7   s&     r>   Ú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_impl]   s    

rC   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ÚscalarrC   rA   r>   Úcast)
rD   rE   r"   Zallow_lhs_ptrZallow_rhs_ptrZarithmetic_checkr4   Z
lhs_sca_tyZ
rhs_sca_tyÚ
ret_sca_tyr   r   r   Úbinary_op_type_checking_impli   s    rK   )ÚinputÚotherr"   r#   c                 C  s°   t | ||ddƒ\} }| jj}|jj}| ¡ r>| ¡ s>||  } }| ¡ r`t | | j|j¡| j¡S | ¡ r‚t | 	| j|j¡| j¡S | 
¡ r¤t | | j|j¡| j¡S ds¬t‚d S ©NTF)rK   rG   rH   rA   r%   r&   Úcreate_addptrÚhandlerB   Zcreate_faddr<   Z
create_addr1   ©rL   rM   r"   Úinput_scalar_tyÚother_scalar_tyr   r   r   Úadd}   s    
rT   c                 C  s”   t | ||ddƒ\} }| jj}| ¡ rDt | | jt||ƒj¡| j¡S | 	¡ rft | 
| j|j¡| j¡S | ¡ rˆt | | j|j¡| j¡S ds�t‚d S rN   )rK   rG   rH   rA   r%   r&   rO   rP   ÚminusrB   Zcreate_fsubr<   Z
create_subr1   ©rL   rM   r"   Ú	scalar_tyr   r   r   Úsub“   s    ÿrX   c                 C  sh   t | ||ƒ\} }| jj}| ¡ r:t | | j|j¡| j¡S | ¡ r\t | 	| j|j¡| j¡S dsdt
‚d S r-   )rK   rG   rH   rB   r%   r&   Zcreate_fmulrP   r<   Z
create_mulr1   rV   r   r   r   Úmul¥   s    rY   c                 C  sì   t | ||ddddƒ\} }| jj}|jj}| ¡ rF| ¡ rFt|||ƒ}nŒ| ¡ rd| ¡ rdt| ||ƒ} nn| ¡ r’| ¡ r’t| tj|ƒ} t|tj|ƒ}n@| ¡ rÊ| ¡ rÊ|j|jkr¼t|||ƒ}qÒt| ||ƒ} ndsÒt	‚t 
| | j|j¡| j¡S ©NFT)rK   rG   rH   rB   r<   rI   r%   r7   Zfp_mantissa_widthr1   r&   Úcreate_fdivrP   rQ   r   r   r   Útruediv³   s     r\   c                 C  s¢   t | ||ddddƒ\} }| jj}|jj}| ¡ r–| ¡ r–t||ƒ}t| ||ƒ} t|||ƒ}| ¡ r|t | 	| j
|j
¡| j¡S t | | j
|j
¡| j¡S dsžt‚d S rZ   )rK   rG   rH   r<   r2   rI   Úis_int_signedr%   r&   Zcreate_sdivrP   Zcreate_udivr1   )rL   rM   r"   rR   rS   Úret_tyr   r   r   ÚfloordivÏ   s    
r_   )rL   rM   Ú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)	rG   rH   rB   r=   rK   r[   rP   r%   r&   )rL   rM   r`   r"   rR   rS   Úretr   r   r   Úfdivà   s    rb   c              	   C  sÔ   t | ||ddddƒ\} }| jj}|jj}| ¡ rXt| ttt| |d|ƒ|ƒ||ƒ|ƒ}|S | ¡ rÈ|j	|j	krŒt
d| ¡  d | ¡  d ƒ‚| ¡ r®t | | j|j¡| j¡S t | | j|j¡| j¡S dsÐt‚d S )NFTzCannot mod z by r5   )rK   rG   rH   rB   rX   rY   Úfloorrb   r<   r/   r=   r   r]   r%   r&   Zcreate_sremrP   Zcreate_uremr1   )rL   rM   r"   rW   rS   ra   r   r   r   Úmodí   s$     ÿþ rd   c                 C  sz   t | ||dddƒ\} }| jj}|jj}| ¡ r6| ¡ s@t||ƒ‚t||ƒ}||kr^t| ||ƒ} ||krrt|||ƒ}| |fS r-   )rK   rG   rH   r<   r   r2   rI   )rL   rM   r"   Úinput_sca_tyZother_sca_tyrJ   r   r   r   Úbitwise_op_type_checking_impl  s    

rf   c                 C  s*   t | ||ƒ\} }t | | j|j¡| j¡S r$   )rf   r%   r&   Z
create_andrP   rG   ©rL   rM   r"   r   r   r   Úand_  s    rh   c                 C  s*   t | ||ƒ\} }t | | j|j¡| j¡S r$   )rf   r%   r&   Z	create_orrP   rG   rg   r   r   r   Úor_"  s    ri   c                 C  s*   t | ||ƒ\} }t | | j|j¡| j¡S r$   )rf   r%   r&   Z
create_xorrP   rG   rg   r   r   r   Úxor_)  s    rj   c                 C  sD   | j  ¡ st| t d¡|ƒ} |j  ¡ s8t|t d¡|ƒ}t| ||ƒS ©NÚint1)rG   Úis_int1Úbitcastr%   r0   rh   rg   r   r   r   Úlogical_and0  s
    

ro   c                 C  sD   | j  ¡ st| t d¡|ƒ} |j  ¡ s8t|t d¡|ƒ}t| ||ƒS rk   )rG   rm   rn   r%   r0   ri   rg   r   r   r   Ú
logical_or8  s
    

rp   ©rL   r"   c                 C  s&   | j  ¡ st| t d¡|ƒ} t| |ƒS rk   )rG   rm   rn   r%   r0   Úinvertrq   r   r   r   Únot_@  s    
rs   c                 C  s*   t | ||ƒ\} }t | | j|j¡| j¡S r$   )rf   r%   r&   Zcreate_lshrrP   rG   rg   r   r   r   ÚlshrF  s    rt   c                 C  s*   t | ||ƒ\} }t | | j|j¡| j¡S r$   )rf   r%   r&   Zcreate_ashrrP   rG   rg   r   r   r   ÚashrM  s    ru   c                 C  s*   t | ||ƒ\} }t | | j|j¡| j¡S r$   )rf   r%   r&   Z
create_shlrP   rG   rg   r   r   r   ÚshlT  s    rv   )rL   r#   c                 C  s   | S r$   r   )rL   r   r   r   Úplus_  s    rw   )rL   r"   r#   c                 C  sH   | j j}| ¡ r$td| ¡  d ƒ‚t | | |¡¡|¡}t	|| |ƒS )Nz$wrong type argument to unary minus (ú))
rG   rH   rA   r=   r   r%   r&   Úget_null_valueÚto_irrX   )rL   r"   re   Ú_0r   r   r   rU   c  s
    rU   c                 C  sP   | j j}| ¡ s| ¡ r,td| ¡  d ƒ‚t | | 	|¡¡|¡}t
| ||ƒS )Nz%wrong type argument to unary invert (rx   )rG   rH   rA   rB   r=   r   r%   r&   Zget_all_ones_valuerz   rj   )rL   r"   re   Z_1r   r   r   rr   l  s
    rr   ztl.block_type)Úvr#   c                 C  s&   | j  ¡ stjS | j j}t tj|¡S r$   )rG   Úis_blockr%   rl   ÚshapeÚ
block_type)r|   r~   r   r   r   Ú
_bool_likex  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sŒt‚d S r-   )rK   rG   rH   rB   r%   r&   Zcreate_fcmpOGTrP   r€   r<   r]   Zcreate_icmpSGTZcreate_icmpUGTr1   rV   r   r   r   Úgreater_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sŒt‚d S r-   )rK   rG   rH   rB   r%   r&   Zcreate_fcmpOGErP   r€   r<   r]   Zcreate_icmpSGEZcreate_icmpUGEr1   rV   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sŒt‚d S r-   )rK   rG   rH   rB   r%   r&   Zcreate_fcmpOLTrP   r€   r<   r]   Zcreate_icmpSLTZcreate_icmpULTr1   rV   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sŒt‚d S r-   )rK   rG   rH   rB   r%   r&   Zcreate_fcmpOLErP   r€   r<   r]   Zcreate_icmpSLEZcreate_icmpULEr1   rV   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sht‚d S r-   )rK   rG   rH   rB   r%   r&   Zcreate_fcmpOEQrP   r€   r<   Zcreate_icmpEQr1   rV   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sht‚d S r-   )rK   rG   rH   rB   r%   r&   Zcreate_fcmpUNErP   r€   r<   Zcreate_icmpNEr1   rV   r   r   r   Ú	not_equalÑ  s    r†   )ÚstartÚendr"   r#   c                 C  s€   t | tƒrt |tƒstdƒ‚t| d? ƒ}t|d? ƒ}|s<|rDtdƒ‚|| krTtdƒ‚||  g}t tj|¡}t | | |¡|¡S )Nz/arange's arguments must be of type tl.constexpré    zarange must fit in int32z=arange's end argument must be greater than the start argument)	Ú
isinstancer    r=   r3   r%   r   r'   r&   Zcreate_make_range)r‡   rˆ   r"   Zis_start_int64Zis_end_int64r~   r^   r   r   r   Úarangeã  s    
r‹   z	List[int])r~   r0   r"   r#   c                 C  s¸   t |tjƒrP|jjdks tdƒ‚t|||ƒ}t |j| ¡}t | 	|j
| ¡|¡S |dkrj| | |¡¡}nt|d|j› �ƒ}||ƒ}|d kr”tdƒ‚t || ¡}t | 	|| ¡|¡S d S )Nr   zonly accepts size-1 tensorr   Úget_z2dtype must be specified when value is not a tensor)rŠ   r%   r&   ÚnumelÚvaluer1   rI   r   r0   Úcreate_splatrP   ry   rz   ÚgetattrÚnamer=   )r~   rŽ   r0   r"   r^   Zget_value_fnr   r   r   Úfullò  s    r’   )rL   Ú	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)	rG   r�   r=   r%   r   rH   r&   Zcreate_viewrP   )rL   r“   r"   r�   Úsr^   r   r   r   Úview
  s    
r•   c                 C  s   t dƒ‚d S )Nzž`reshape` is not supported yet. Please use `view` instead if applicable. Note that view may reorder elements in an implementation- and context- dependent way.)r=   )rL   r“   r"   r   r   r   Úreshape  s    r–   )rL   r!   r"   r#   c                 C  s>   t | jjƒ}| |d¡ t | jj|¡}t | | j	|¡|¡S )Nr   )
ÚlistrG   r~   Úinsertr%   r   rH   r&   Úcreate_expand_dimsrP   )rL   r!   r"   r“   r^   r   r   r   Úexpand_dims   s    rš   )rD   rE   Úcan_reorderr"   r#   c                 C  sX   |st dƒ‚t| jƒdkst ‚t | jj| jd |jd  g¡}t | | j	|j	¡|¡S )Nz;current implementation of `cat` always may reorder elementsr   r   )
r1   Úlenr~   r%   r   rG   rH   r&   Z
create_catrP   )rD   rE   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 )Nr
   z!Only 2D tensors can be transposedr   r   )
rœ   r~   r=   r%   r   rG   rH   r&   Zcreate_transrP   )rL   r"   r�   r   r   r   Útrans.  s     rŸ   )rL   r~   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 ú: )rG   r}   r%   r   r&   r�   rP   Úget_block_shapesrœ   r=   Ú	enumeraterH   Úcreate_broadcast)rL   r~   r"   r^   Z	src_shapeÚiÚitemr   r   r   Úbroadcast_impl_shape5  s    

,r¦   c              	   C  sp  | j }|j }| ¡ rJ| ¡ sJt |j|j¡}t | |j| 	¡ ¡|¡}�n| ¡ sˆ| ¡ rˆt |j|j¡}t | | j| 	¡ ¡|¡} �nà| ¡ �rh| ¡ �rh| 	¡ }| 	¡ }t
|ƒt
|ƒk �rtt
|ƒt
|ƒƒD ]8}t | | jd¡t |jdg| ¡¡} | j }| 	¡ }qÐn`t
|ƒt
|ƒk �rltt
|ƒt
|ƒƒD ]:}t | |jd¡t |jdg| ¡¡}|j }| 	¡ }�q0t
|ƒt
|ƒk�s‚t‚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   r   z?Cannot make_shape_compatible: incompatible dimensions at index r    r   )rG   r}   r%   r   rH   r~   r&   r�   rP   r¡   rœ   Úranger™   r1   r¢   Úappendr=   Ústrr£   )rD   rE   r"   Zlhs_tyZrhs_tyZ	lhs_shapeZ	rhs_shapeÚdimÚ	ret_shaper¤   ÚleftÚrightr^   r   r   r   rF   I  sb    &&


ÿÿÿÿÿ


rF   )rL   Údst_tyr"   r#   c                 C  s¢   | j }| ¡ r"t |j| j  ¡ ¡}||kr.| S |j}|j}| ¡ sJ| ¡ rVt| ||ƒS |j}|j}||kr†t	dt
|ƒ d t
|ƒ ƒ‚t | | j| |¡¡|¡S )Nz!Cannot bitcast data-type of size z to data-type of size )rG   r}   r%   r   rH   r¡   rA   rI   Úprimitive_bitwidthr=   r©   r&   Úcreate_bitcastrP   rz   )rL   r®   r"   Úsrc_tyÚ
src_sca_tyÚ
dst_sca_tyZsrc_bitsZdst_bitsr   r   r   rn   „  s$    ÿÿrn   c                 C  s
   t | tƒS r$   ©rŠ   r    )Úarchr   r   r   Ú_is_cuda›  s    r¶   c                 C  sä  | j }t|tjƒr|j}| ¡ r4t |j| j  ¡ ¡}||kr@| S |j}|j}t	|j
ƒr||j
dk r|| ¡ sp| ¡ r|t dt¡ | ¡ rŒ| ¡ sœ| ¡ r¸| ¡ r¸t | | j| |¡¡|¡S | ¡ rÈ| ¡ rØ| ¡ rî| ¡ sîtt| tj|ƒ||ƒS | ¡ �o| ¡ �o|j|jk}|�r0t | | j| |¡¡|¡S | ¡ �oN| ¡ �oN|j|jk }|�rrt | | j| |¡¡|¡S | ¡ �r
| ¡ �r
|j|jk�s¢|j|jk�r
|  ¡ �o´| !¡  }| !¡ �rì| j" |¡}	t | #|	¡| j"¡}
t$| |
|ƒS t | %| j| |¡|¡|¡S | &¡ �r–| ¡ �r–| !¡ �rT| j" |¡}	t | #|	¡| j"¡}
t$| |
|ƒS |  ¡ �rzt | '| j| |¡¡|¡S t | (| j| |¡¡|¡S | ¡ �rö| &¡ �rö| !¡ �s¾|  ¡ �sÚt | )| j| |¡¡|¡S t | *| j| |¡¡|¡S | +¡ �rf| ¡ �rf|j}|dk�r6t | ,| j| |¡¡|¡S |dk�rft$t| tj-|ƒt | .d¡tj-¡|ƒS | ¡ �r–| +¡ �r–t | /| j| |¡¡|¡S | +¡ �rÆ| +¡ �rÆt | 0| j| |¡¡|¡S d�sàt1d| › d|› �ƒ‚d S )	NéY   zUStandard tl.float8e4 format will be deprecated on SM < 89. Please use tl.float8e4b15.é@   r   r   Fzcannot cast z to )2rG   rŠ   r%   Ú	constexprrŽ   r}   r   rH   r¡   r¶   rµ   Zis_fp8e4ÚwarningsÚwarnÚDeprecationWarningZis_fp8rB   r&   Zcreate_fp_to_fprP   rz   r8   r6   r:   rI   r7   r¯   Zcreate_fp_truncZcreate_fp_extr<   r.   r/   r]   Úis_boolr0   ry   r†   Úcreate_int_castZis_standard_floatingZcreate_fp_to_siZcreate_fp_to_uiZcreate_ui_to_fpZcreate_si_to_fprA   Zcreate_ptr_to_intÚint64Ú	get_int64Zcreate_int_to_ptrr°   r1   )rL   r®   r"   r±   r²   r³   Ztruncate_fpZext_fpZsign_extendÚtyr{   Zbitwidthr   r   r   rI   Ÿ  sâ    ÿÿÿÿÿÿÿÿ
ÿ
þÿþ
ÿ
þÿþ
ÿ
ÿ
 ÿþ

ÿþÿþÿþÿþ
ÿ
þrI   c                 C  sD   t jj}| r@| dkrt jj}n"| dkr0t jj}ntd| › d�ƒ‚|S )Nz.caú.cgúCache modifier ú not supported)r   ÚCACHE_MODIFIERÚNONEÚCAÚCGr=   ©Úcache_modifierÚcacher   r   r   Ú_str_to_load_cache_modifier  s    

rÌ   c                 C  sh   t jj}| rd| dkrt jj}nF| dkr0t jj}n4| dkrBt jj}n"| dkrTt jj}ntd| › d�ƒ‚|S )Nz.wbrÂ   z.csz.wtrÃ   rÄ   )r   rÅ   rÆ   ZWBrÈ   ÚCSZWTr=   rÉ   r   r   r   Ú_str_to_store_cache_modifier!  s    



rÎ   c                 C  sD   t jj}| r@| dkrt jj}n"| dkr0t jj}ntd| › d�ƒ‚|S )NZ
evict_lastZevict_firstzEviction policy rÄ   )r   ZEVICTION_POLICYZNORMALZ
EVICT_LASTZEVICT_FIRSTr=   )Úeviction_policyÚevictionr   r   r   Ú_str_to_eviction_policy1  s    

rÑ   c                 C  s@   d }| r<| dkrt jj}n"| dkr,t jj}ntd| › d�ƒ‚|S )NÚzeroÚnanzPadding option rÄ   )r   ÚPADDING_OPTIONZPAD_ZEROÚPAD_NANr=   )Úpadding_optionÚpaddingr   r   r   Ú_str_to_padding_option=  s    

rØ   c                 C  sh   t jj}| rd| dkrt jj}nF| dkr0t jj}n4| dkrBt jj}n"| dkrTt jj}ntd| › d�ƒ‚|S )NÚacquireÚreleaseZacq_relÚrelaxedzMemory semantic rÄ   )r   ZMEM_SEMANTICZACQUIRE_RELEASEZACQUIREZRELEASEZRELAXEDr=   )Z
sem_optionÚsemr   r   r   Ú_str_to_semI  s    



rÝ   c                 C  sŽ   | rˆt | dƒs| g} dd„ | D ƒ} | D ],}t|tƒrNd|  krLt|ƒk s&n t‚q&t| ƒdksdt‚t| ƒtt| ƒƒks€tdƒ‚t| ƒS tƒ S )NÚ__iter__c                 S  s"   g | ]}t |tjƒr|jn|‘qS r   ©rŠ   r%   r¹   rŽ   ©Ú.0Úelemr   r   r   Ú
<listcomp>]  s     z0_canonicalize_boundary_check.<locals>.<listcomp>r   z'Duplicate dimension in `boundary_check`)ÚhasattrrŠ   r    rœ   r1   ÚsetÚsortedÚtuple)Úboundary_checkÚblock_shaperª   r   r   r   Ú_canonicalize_boundary_checkY  s    
*rê   c	              
   C  s|   |s|rt dƒ‚| jjj}	|	tjks,tdƒ‚|	 ¡ rH|tjj	krHt dƒ‚| jj}
t
||
 ¡ ƒ}t | | j|||||¡|
¡S )NúK`mask` and `other` arguments cannot be specified for loading block pointersú3`tl.int1` should be rewrited in `tl.make_block_ptr`z@Padding option `nan` is not supported for integer block pointers)r=   rG   Ú
element_tyr%   rl   r1   r<   r   rÔ   rÕ   rê   r¡   r&   Zcreate_tensor_pointer_loadrP   )ÚptrÚmaskrM   rè   r×   rË   rÐ   Úis_volatiler"   Úelt_tyr®   r   r   r   Ú_load_block_pointerf  s    
ÿÿrò   c	              
   C  sn  | j j ¡ s"td| j  ¡ › d�ƒ‚|s2|r2tdƒ‚|s:|rBtdƒ‚| j  ¡ sx|rb|j  ¡ rbtdƒ‚|rx|j  ¡ rxtdƒ‚| j  ¡ r®|r˜t|| j  ¡ |ƒ}|r®t|| j  ¡ |ƒ}| j j}	|	j}
|
t	j
kræt	j}
t	 |
|	j¡}	t| |	|ƒ} |röt||
|ƒ}| j  ¡ �r| j  ¡ }t	 |
|¡}n|
}|�s>t	 | | j|||¡|¡S t	 | | j|j|�rZ|jnd |||¡|¡S d S )NúUnsupported ptr type z in `tl.load`z)`other` cannot be provided without `mask`zä`padding_option` or `boundary_check` argument is not supported for loading a tensor ofpointers or loading a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadúEMask argument cannot be block type if pointer argument is not a blockzFOther argument cannot be block type if pointer argument is not a block)rG   rH   rA   r=   r   r}   r¦   r¡   rí   r%   rl   Úint8Úpointer_typeÚaddress_spacerI   r   r&   Zcreate_loadrP   Zcreate_masked_load)rî   rï   rM   rè   r×   rË   rÐ   rð   r"   Úptr_tyrñ   r~   r®   r   r   r   Ú_load_legacy|  sH    



  ÿÿrù   zOptional[tl.tensor]r©   )	rî   rï   rM   rÖ   rÊ   rÏ   rð   r"   r#   c	              
   C  sb   t |ƒ}	t|ƒ}
t|ƒ}| j ¡ rF| jj ¡ rFt| |||||	|
||ƒ	S t| |||||	|
||ƒ	S d S r$   )	rÌ   rÑ   rØ   rG   rA   rí   r}   rò   rù   )rî   rï   rM   rè   rÖ   rÊ   rÏ   rð   r"   rË   rÐ   r×   r   r   r   Úloadµ  s    
rú   c           	   	   C  s¶   |rt dƒ‚| jj ¡ }|j ¡ s.t|||ƒ}|j ¡ s@tdƒ‚||j ¡ ksVtdƒ‚| jjj|jjksptdƒ‚| jjj}|tjksŒtdƒ‚t	||ƒ}t 
| | j|j|||¡tj¡S )Nrë   z-Value argument must be block type or a scalarz$Block shape and value shape mismatchz2Block element type and value element type mismatchrì   )r=   rG   rí   r¡   r}   r¦   r1   r%   rl   rê   r&   Zcreate_tensor_pointer_storerP   Úvoid)	rî   Úvalrï   rè   rË   rÐ   r"   ré   rñ   r   r   r   Ú_store_block_pointerË  s    


ÿrý   c           	   	   C  s0  | j j ¡ s"td| j  ¡ › d�ƒ‚|r.tdƒ‚| j  ¡ s`|j  ¡ rJtdƒ‚|r`|j  ¡ r`tdƒ‚| j  ¡ r’t|| j  ¡ |ƒ}|r’t|| j  ¡ |ƒ}| j j}|j}|t	j
krÊt	j}t	 ||j¡}t| ||ƒ} t|||ƒ}|søt	 | | j|j||¡t	j¡S |j j ¡ �stdƒ‚t	 | | j|j|j||¡t	j¡S )Nró   z in `tl.store`zÑ`boundary_check` argument is not supported for storing a tensor of pointers or storing a scalar. Because the compiler does not know the boundary; please use block pointers (defined by `make_block_ptr`) insteadzFValue argument cannot be block type if pointer argument is not a blockrô   z"Mask must have boolean scalar type)rG   rH   rA   r=   r   r}   r¦   r¡   rí   r%   rl   rõ   rö   r÷   rI   r&   Zcreate_storerP   rû   r½   Zcreate_masked_store)	rî   rü   rï   rè   rË   rÐ   r"   rø   rñ   r   r   r   Ú_store_legacyä  s2    



rþ   )rî   rü   rï   rÊ   rÏ   r"   r#   c           	      C  sR   t |ƒ}t|ƒ}| j ¡ r:| jj ¡ r:t| ||||||ƒS t| ||||||ƒS d S r$   )rÎ   rÑ   rG   rA   rí   r}   rý   rþ   )	rî   rü   rï   rè   rÊ   rÏ   r"   rË   rÐ   r   r   r   Ústore  s
    rÿ   )rî   Úcmprü   rÜ   r"   r#   c                 C  sD   t |ƒ}| jjj}|jdkr$tdƒ‚t | | j	|j	|j	|¡|j¡S )N)é   r‰   r¸   z9atomic_cas only supports elements with width {16, 32, 64})
rÝ   rG   rH   rí   r¯   r=   r%   r&   Zcreate_atomic_casrP   )rî   r   rü   rÜ   r"   rí   r   r   r   Ú
atomic_cas(  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krJ|dkrJtd| d ƒ‚|tjtj	tj
tjfkrztd| d t|ƒ ƒ‚| j  ¡ r°|ršt|| j  ¡ |ƒ}|r°t|| j  ¡ |ƒ}t|| j jj|ƒ}|�s| d¡}tj}| j  ¡ �r
| || j  ¡ ¡}t tj| j  ¡ ¡}t ||¡}| ||fS )Nz)Pointer argument of store instruction is rT   Zatomic_z does not support fp16z does not support T)rG   rH   rA   r=   r   rí   r%   r9   rl   rõ   Zint16r;   r©   r}   r¦   r¡   rI   Zget_int1r�   r   r&   )rî   rü   rï   r  r"   rí   Zmask_irZmask_tyr   r   r   Úatom_red_typechecking_impl4  s*    


r  )rî   rü   rï   rÜ   r"   r#   c                 C  sD  t | ||d|ƒ\} }}t|ƒ}|jj}| ¡ r‚| ¡ r\t | 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  rÝ   rG   rH   r<   r]   r%   r&   Úcreate_atomic_rmwr   Ú	ATOMIC_OPÚMAXrP   ÚUMAXrn   r'   rö   r‚   Úget_fp32rƒ   rh   ÚUMINÚwhere©rî   rü   rï   rÜ   r"   Úsca_tyZi_valZi_ptrÚposÚnegZpos_retZneg_retr   r   r   Ú
atomic_maxP  s8    üûüû	..r  c                 C  sD  t | ||d|ƒ\} }}t|ƒ}|jj}| ¡ r‚| ¡ r\t | 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  rÝ   rG   rH   r<   r]   r%   r&   r  r   r  ZMINrP   r  rn   r'   rö   r‚   r
  rƒ   rh   r	  r  r  r   r   r   Ú
atomic_mint  sT    üûüû	üûüûr  c              	   C  s`   t | ||d|ƒ\} }}t|ƒ}|jj}| ¡ r6tjjntjj}t	 
| || j|j|j|¡|j¡S )NrT   )r  rÝ   rG   rH   rB   r   r  ZFADDZADDr%   r&   r  rP   )rî   rü   rï   rÜ   r"   r  r  r   r   r   Ú
atomic_add¢  s
    r  c              	   C  sD   t | ||d|ƒ\} }}t|ƒ}t | tjj| j|j|j|¡|j	¡S )NÚand)
r  rÝ   r%   r&   r  r   r  ÚANDrP   rG   ©rî   rü   rï   rÜ   r"   r   r   r   Ú
atomic_and®  s    r  c              	   C  sD   t | ||d|ƒ\} }}t|ƒ}t | tjj| j|j|j|¡|j	¡S )NÚor)
r  rÝ   r%   r&   r  r   r  ÚORrP   rG   r  r   r   r   Ú	atomic_or¸  s    r  c              	   C  sD   t | ||d|ƒ\} }}t|ƒ}t | tjj| j|j|j|¡|j	¡S )NÚxor)
r  rÝ   r%   r&   r  r   r  ZXORrP   rG   r  r   r   r   Ú
atomic_xorÂ  s    r  c              	   C  sD   t | ||d|ƒ\} }}t|ƒ}t | tjj| j|j|j|¡|j	¡S )NZxchg)
r  rÝ   r%   r&   r  r   r  ZXCHGrP   rG   r  r   r   r   Úatomic_xchgÌ  s    r  )rD   rE   Ú
allow_tf32Ú	out_dtyper"   r#   c           
   
   C  s  | j  ¡ r|j  ¡ st‚| j|jks>td| j› d|j› d�ƒ‚t| jƒdks^td| j› d�ƒ‚t|jƒdks~td|j› d�ƒ‚| jd j|jd	 jksÌtd| j› d
|j› d| jd j› d|jd	 j› d�	ƒ‚| jd	 jdkrþ| jd jdkrþ|jd jdk�std| j› d|j› d�ƒ‚| j j ¡ �rj| j jt	j
k�s>tdƒ‚| jd jdk�sXtdƒ‚| d	¡}t	j}nP| j j ¡ �s†| j j ¡ �r˜| d	¡}t	j}n"| ¡ �r¬| d	¡n| d	¡}|}| j jd	 }|j jd }| |||g¡}t	 |||g¡}	t	 | | j|j||¡|	¡S )NzFirst input (z) and second input (z) must have the same dtype!r
   zFirst input shape (z) is not two dimensional!zSecond input shape (r   r   z) and second input shape z= are not compatible for matmul (second index of first shape (z0) must be equal to first index of second shape (rx   r  z&All values in both first input shape (z) and second input shape (z) must be >= 16!zonly int8 supported!r‰   zsmall blocks not supported!)rG   r}   r1   r0   rœ   r~   rŽ   rH   r<   r%   rõ   Ú	get_int32r'   r6   r:   r
  r7   r8   Zget_fp16r�   r   r&   Z
create_dotrP   )
rD   rE   r  r   r"   r{   Zret_scalar_tyÚMÚNr^   r   r   r   ÚdotÚ  s6    &  N ÿþ

ÿ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)
rI   r%   rl   rG   r}   rF   rK   r&   Zcreate_selectrP   )r%  r&  r'  r"   Ú_r^   r   r   r   r     s    

r  zSequence[tl.tensor]zTuple[tl.tensor, ...])Úinputsr!   r"   r#   c           	        sÞ   ˆ d krNg }t tˆƒƒD ](}ˆ| jjg}| tˆ| ||ƒ¡ qt|ƒ‰d‰ ˆd jj}‡ fdd„t	|ƒD ƒ‰ˆD ]}|jj|kstt
‚qt‡fdd„‰| dd„ ˆD ƒˆ ¡‰|ˆƒ ˆ ¡  t‡‡‡fdd„t tˆƒƒD ƒƒS )	Nr   c                   s   g | ]\}}|ˆ kr|‘qS r   r   )rá   r¤   r”   )r!   r   r   rã   !  s      zreduction.<locals>.<listcomp>c                   s"   ˆ rt  |ˆ ¡}n|}t  | |¡S r$   ©r%   r   r&   ©r&  rW   Zres_ty)r«   r   r   Úwrap_tensor%  s    zreduction.<locals>.wrap_tensorc                 S  s   g | ]
}|j ‘qS r   ©rP   ©rá   Útr   r   r   rã   -  s     c                 3  s&   | ]}ˆˆ  |¡ˆ | jjƒV  qd S r$   ©Z
get_resultrG   rH   ©rá   r¤   )r)  Ú	reduce_opr,  r   r   Ú	<genexpr>1  s   ÿzreduction.<locals>.<genexpr>)r§   rœ   r�   rŽ   r¨   r•   rç   rG   r~   r¢   r1   Zcreate_reduceÚverify)	r)  r!   Úregion_builder_fnr"   Z
new_inputsr¤   Z	new_shaper~   r/  r   )r!   r)  r2  r«   r,  r   Ú	reduction  s$    
þr6  c                   st   t ˆ ƒdkrtdƒ‚ˆ d jj‰‡fdd„‰| dd„ ˆ D ƒ|¡‰|ˆƒ ˆ ¡  t‡ ‡‡fdd	„tt ˆ ƒƒD ƒƒS )
Nr   z7Current implementation only support single tensor inputr   c                   s   t  |ˆ ¡}t  | |¡S r$   r*  r+  )r~   r   r   r,  C  s    z%associative_scan.<locals>.wrap_tensorc                 S  s   g | ]
}|j ‘qS r   r-  r.  r   r   r   rã   G  s     z$associative_scan.<locals>.<listcomp>c                 3  s&   | ]}ˆˆ  |¡ˆ | jjƒV  qd S r$   r0  r1  )r)  Úscan_opr,  r   r   r3  K  s   ÿz#associative_scan.<locals>.<genexpr>)rœ   r=   rG   r~   Zcreate_scanr4  rç   r§   )r)  r!   r5  r"   r   )r)  r7  r~   r,  r   Úassociative_scan<  s    
þr8  z	List[str])Údtypesr#   c                   s   ‡ fdd„}|S )ai  
    We following libdevice's convention to check accepted data types for math functions.
    It is not a good practice to support all data types as accelerators/GPUs don't support
    many float16 and bfloat16 math operations.
    We should let the users know that they are using and invoke explicit cast to convert
    the data type to the supported one.
    c                   s   t ˆ ƒ‡‡ fdd„ƒ}|S )Nc                    sZ   t | ƒt | ¡ ƒ }dd„ |D ƒD ],}|jjjˆ kr"tdˆ › d|jjj› �ƒ‚q"ˆ| |ŽS )Nc                 S  s   g | ]}t |tjƒr|‘qS r   )rŠ   r%   r&   )rá   Úar   r   r   rã   b  s      z@_check_dtype.<locals>.wrapper.<locals>.check.<locals>.<listcomp>zExpected dtype z	 but got )r—   ÚvaluesrG   rH   r‘   r=   )ÚargsÚkwargsZall_argsÚarg)r9  Úfnr   r   Úcheck^  s
    z,_check_dtype.<locals>.wrapper.<locals>.checkr   )r?  r@  ©r9  )r?  r   Úwrapper]  s    z_check_dtype.<locals>.wrapperr   )r9  rB  r   rA  r   Ú_check_dtypeU  s    rC  )r&  r'  r"   r#   c                 C  s,   t | ||ƒ\} }ddlm} |j| ||d�S ©Nr   )Úmath)Z_builder)rK   Ú rE  Zmulhi)r&  r'  r"   rE  r   r   r   Úumulhik  s    rG  Zfp32Zfp64rA  )r&  r"   r#   c                 C  s   ddl m} |j| |d�S rD  )rF  rE  rc   )r&  r"   rE  r   r   r   rc   r  s    rc   c                 C  s   t  | | j¡| j¡S r$   )r%   r&   Z
create_exprP   rG   ©r&  r"   r   r   r   Úexpy  s    rI  c                 C  s   t  | | j¡| j¡S r$   )r%   r&   Z
create_logrP   rG   rH  r   r   r   Úlog~  s    rJ  c                 C  s   t  | | j¡| j¡S r$   )r%   r&   Z
create_cosrP   rG   rH  r   r   r   Úcosƒ  s    rK  c                 C  s   t  | | j¡| j¡S r$   )r%   r&   Z
create_sinrP   rG   rH  r   r   r   Úsinˆ  s    rL  c                 C  s   t  | | j¡| j¡S r$   )r%   r&   Zcreate_sqrtrP   rG   rH  r   r   r   Úsqrt�  s    rM  c                 C  sd   | j }| ¡ r$t | | j¡| j¡S | ¡ rBt | | j¡| j¡S | 	¡ rN| S ds`t
d|› �ƒ‚d S )NFzUnexpected dtype )r0   rB   r%   r&   Zcreate_fabsrP   rG   r]   Zcreate_iabsZis_int_unsignedr1   )r&  r"   r0   r   r   r   Úabs’  s    rN  )r&  r;  r#   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œ   r~   r=   rP   Zset_attrr   Z	make_attrZget_context©r&  r;  r   r   r   Úmultiple_of¡  s    rQ  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.contiguityrO  rP  r   r   r   Úmax_contiguous¨  s    rR  c                 C  s:   t | jƒt |ƒkrtdƒ‚| j dt || j ¡ ¡¡ | S )NzCShape of input to max_constancy does not match the length of valuesztt.constancyrO  rP  r   r   r   Úmax_constancy¯  s    rS  )r"   r#   c                 C  s   t  |  ¡ t j¡S r$   )r%   r&   Zcreate_barrierrû   )r"   r   r   r   Údebug_barrier¶  s    rT  zList[tl.tensor])Úprefixr<  r"   r#   c                 C  s0   g }|D ]}|  |j¡ qt | | |¡tj¡S r$   )r¨   rP   r%   r&   Zcreate_printrû   )rU  r<  r"   Únew_argsr>  r   r   r   Údevice_printº  s    rW  )ÚcondÚmsgÚ	file_nameÚlinenor"   r#   c              	   C  sP   | j }| ¡ s2t |jd¡}t | | jd¡|¡} t | | j||||¡tj	¡S )N)r   )
rG   r}   r%   r   rH   r&   r�   rP   Zcreate_assertrû   )rX  rY  rZ  Ú	func_namer[  r"   Zcond_tyr   r   r   Údevice_assertÁ  s
    r]  c                 C  sÊ   t |tjƒr(|r|  |j¡S |  |j¡S t |tjƒr°|jjdksHtdƒ‚|j	 
¡ sZtdƒ‚|j	tjkr„|r„|  |j|  ¡ |j	 ¡ ¡S |j	tjkrª|  |j|  ¡ |j	 ¡ ¡S |jS dsÆtdt|ƒ› �ƒ‚d S )Nr   z*Expected a scalar in shape/strides/offsetsz8Expected an integer scalar type in shape/strides/offsetsFz3Unsupported element type in shape/strides/offsets: )rŠ   r%   r¹   rÀ   rŽ   r!  r&   r�   r1   r0   r<   r¿   r¾   rP   Zget_int64_tyr]   r'   Zget_int32_tyrG   )r"   râ   Úrequire_i64r   r   r   Ú_convert_elem_to_ir_valueÉ  s    r_  c                   s,   t |dƒr‡ ‡fdd„|D ƒS tˆ |ˆƒgS )NrÞ   c                   s   g | ]}t ˆ |ˆƒ‘qS r   )r_  rà   ©r"   r^  r   r   rã   Ù  s     z)_convert_to_ir_values.<locals>.<listcomp>)rä   r_  )r"   Ú	list_liker^  r   r`  r   Ú_convert_to_ir_values×  s    
rb  )Úbaser"   r#   c              	     s:  t ||ƒ}t ||ƒ}t ||dd�}| j ¡ r8| jj ¡ r@tdƒ‚| jjtjkrht| t 	tj
| jj¡|ƒ} tˆ dƒsxˆ g‰ dd„ ˆ D ƒ‰ tdd„ ˆ D ƒƒs tdƒ‚t|dƒs°|g}d	d„ |D ƒ}t|ƒttt|ƒƒƒksÞtd
ƒ‚t‡ fdd„||||fD ƒƒ�stdƒ‚| | j|||ˆ |¡}t |t 	t | jjˆ ¡¡¡S )NF©r^  zMExpected `base` to be a pointer type (but not a block pointer type or others)rÞ   c                 S  s"   g | ]}t |tjƒr|jn|‘qS r   rß   rà   r   r   r   rã   ï  s     z"make_block_ptr.<locals>.<listcomp>c                 S  s.   g | ]&}t |tƒo(d |  ko$dk n  ‘qS )i   €l        r´   rà   r   r   r   rã   ð  s     zGExpected a list of constant integers (`int32_t` range) in `block_shape`c                 S  s"   g | ]}t |tjƒr|jn|‘qS r   rß   rà   r   r   r   rã   ö  s     z<Expected a permutation of (0, 1, ..., len(order)-1) in orderc                   s   g | ]}t ˆ ƒt |ƒk‘qS r   )rœ   )rá   ra  ©ré   r   r   rã   ú  s     zBExpected shape/strides/offsets/block_shape to have the same length)rb  rG   rA   rí   r}   r=   r%   rl   rI   rö   rõ   r÷   rä   Úallr1   ræ   r—   r§   rœ   Zcreate_make_block_ptrrP   r&   r   )rc  r~   ÚstridesÚoffsetsré   Úorderr"   rP   r   re  r   Úmake_block_ptrÝ  s,    


ÿ
 "ÿrj  c                 C  s&   t ||dd�}t | | j|¡| j¡S )NFrd  )rb  r%   r&   Zcreate_advancerP   rG   )rc  rh  r"   r   r   r   Úadvance  s    rk  )FFTF)T)lÚ
__future__r   rº   Ú	functoolsr   Útypingr   r   r   r   r	   Z_C.libtriton.tritonr   rF  r   r%   r   Ú	Exceptionr   r)   r*   r2   r>   rC   rK   rT   rX   rY   r\   r_   rb   rd   rf   rh   ri   rj   ro   rp   rs   rt   ru   rv   rw   rU   rr   r€   r�   r‚   rƒ   r„   r…   r†   r‹   r’   r•   r–   rš   rž   rŸ   r¦   rF   rn   r¶   rI   rÌ   rÎ   rÑ   rØ   rÝ   rê   rò   rù   rú   rý   rþ   rÿ   r  r  r  r  r  r  r  r  r  r$  r  r6  r8  rC  rG  rc   rI  rJ  rK  rL  rM  rN  rQ  rR  rS  rT  rW  r]  r_  rb  rj  rk  r   r   r   r   Ú<module>   sÔ   &      ü	;v9,$.


&'
'