o
    ۾i3                     @   s`  d dl Zd dlmZmZmZmZ d dlmZ d dl	m
Z
mZ d dlmZ d dlZd dlZd dlm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!ej"d'd(d)d*d+ Z#ej"d'd(d)d,d- Z$d.d/ Z%d0d1 Z&d2d3 Z'd4d5 Z(d6d7 Z)G d8d9 d9eZ*e+d:kre,  dS dS );    N)unittestCUDATestCaseskip_unless_cc_53skip_on_cudasim)cuda)f2b1)compile_ptx)
from_dtypec                 C   s   || | d< d S Nr    aryabr   r   Y/home/ubuntu/.local/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_operator.pysimple_fp16_div_scalar      r   c                 C   s   || | d< d S r   r   r   r   r   r   simple_fp16add   r   r   c                 C   s   | d  |7  < d S r   r   r   r   r   r   r   simple_fp16_iadd      r   c                 C   s   | d  |8  < d S r   r   r   r   r   r   simple_fp16_isub   r   r   c                 C   s   | d  |9  < d S r   r   r   r   r   r   simple_fp16_imul   r   r   c                 C   s   | d  |  < d S r   r   r   r   r   r   simple_fp16_idiv    r   r   c                 C   s   || | d< d S r   r   r   r   r   r   simple_fp16sub$   r   r   c                 C   s   || | d< d S r   r   r   r   r   r   simple_fp16mul(   r   r   c                 C   s   | | d< d S r   r   r   r   r   r   simple_fp16neg,   s   r   c                 C   s   t || d< d S r   )absr   r   r   r   simple_fp16abs0   r   r   c                 C   s   ||k| d< d S r   r   r   r   r   r   simple_fp16_gt4   r   r    c                 C   s   ||k| d< d S r   r   r   r   r   r   simple_fp16_ge8   r   r!   c                 C   s   ||k | d< d S r   r   r   r   r   r   simple_fp16_lt<   r   r"   c                 C   s   ||k| d< d S r   r   r   r   r   r   simple_fp16_le@   r   r#   c                 C   s   ||k| d< d S r   r   r   r   r   r   simple_fp16_eqD   r   r$   c                 C   s   ||k| d< d S r   r   r   r   r   r   simple_fp16_neH   r   r%   z
b1(f2, f2)T)devicec                 C      | |k S Nr   xyr   r   r   
hlt_func_1L      r,   c                 C   r'   r(   r   r)   r   r   r   
hlt_func_2Q   r-   r.   c                 C   s   t ||o	t||| d< d S r   )r,   r.   rr   r   cr   r   r   test_multiple_hcmp_1V   s   r2   c                 C   s   t ||o||k | d< d S r   r,   r/   r   r   r   test_multiple_hcmp_2[      r4   c                 C   s   t ||o||k| d< d S r   r3   r/   r   r   r   test_multiple_hcmp_3`   r5   r6   c                 C   s   ||k o||k | d< d S r   r   r/   r   r   r   test_multiple_hcmp_4e      r7   c                 C   s   ||k o||k| d< d S r   r   r/   r   r   r   test_multiple_hcmp_5j   r8   r9   c                       s  e Zd Z fddZ	 dd Zdd Zdd Zd	d
 Zdd Zdd Z	e
dd Zeddd Ze
dd Zeddd Ze
dd Ze
dd Zeddd Zeddd Ze
d d! Ze
d"d# Ze
d$d% Ze
d&d' Zedd(d) Zedd*d+ Zedd,d- Z  ZS ).TestOperatorModulec                    s   t    tjd d S r   )supersetUpnprandomseedself	__class__r   r   r<   p   s   
zTestOperatorModule.setUpc                    sT   t j fdd}td}td}| }|d || tj| || d S )Nc                    s   d} | | || | |< d S r   r   )r   r   iopr   r   foox   s   z1TestOperatorModule.operator_template.<locals>.foo   rH   rH   )r   jitr=   onescopytestingassert_equal)rA   rF   rG   r   r   resr   rE   r   operator_templatew   s   

z$TestOperatorModule.operator_templatec                 C      |  tj d S r(   )rP   operatoraddr@   r   r   r   test_add   r   zTestOperatorModule.test_addc                 C   rQ   r(   )rP   rR   subr@   r   r   r   test_sub   r   zTestOperatorModule.test_subc                 C   rQ   r(   )rP   rR   mulr@   r   r   r   test_mul   r   zTestOperatorModule.test_mulc                 C   rQ   r(   )rP   rR   truedivr@   r   r   r   test_truediv   r   zTestOperatorModule.test_truedivc                 C   rQ   r(   )rP   rR   floordivr@   r   r   r   test_floordiv   r   z TestOperatorModule.test_floordivc           
   	   C   s   t tttf}tjtjtjtjf}t	||D ]U\}}| j
|dC td|}tjdtjd}tjdtj}tjdtj}|d ||d |d  |||}	tj||	 W d    n1 sew   Y  qd S )NrE   zvoid(f2[:], f2, f2)rH   dtyperI   r   )r   r   r   r   rR   rS   rU   rW   rY   zipsubTestr   rJ   r=   zerosfloat16r>   astyperM   assert_allclose
rA   	functionsopsfnrF   kernelgotarg1arg2expectedr   r   r   test_fp16_binary   s    
z#TestOperatorModule.test_fp16_binaryz(Compilation unsupported in the simulatorc              	   C   s   t ttf}d}td d  ttf}t||D ])\}}| j|d t||dd\}}| || W d    n1 s9w   Y  qd S N)zadd.f16zsub.f16zmul.f16)instr      cc)r   r   r   r   r_   r`   r	   assertInrA   rf   instrsargsrh   rp   ptx_r   r   r   test_fp16_binary_ptx   s   
z'TestOperatorModule.test_fp16_binary_ptxc              	   C   s  t tttf}tjtjtjtjf}t	j
t	jt	jt	jt	jt	jf}tt|||D ]]\\}}}| j||dH t|}t	jdt	j}t	jdd |}	t	t	j|}
t	jd|
d}|d ||d |	d  |||	}t	j|| W d    n1 sw   Y  q'd S )NrF   tyrH   d   r]   rI   r   )r   r   r   r   rR   rS   rU   rW   rY   r=   int8int16int32int64float32float64	itertoolsproductr_   r`   r   rJ   r>   rc   rb   result_typera   rM   rd   )rA   rf   rg   typesrh   rF   r~   ri   rk   rl   res_tyrj   rm   r   r   r   !test_mixed_fp16_binary_arithmetic   s(   

z4TestOperatorModule.test_mixed_fp16_binary_arithmeticc              	   C   s   t ttf}d}td d  tf}t||D ])\}}| j|d t||dd\}}| || W d    n1 s8w   Y  qd S ro   )r   r   r   r   r_   r`   r	   rv   rw   r   r   r   test_fp16_inplace_binary_ptx   s   
z/TestOperatorModule.test_fp16_inplace_binary_ptxc           	   	   C   s   t tttf}tjtjtjtjf}t	||D ]N\}}| j
|d< td|}tjdtj}| }tjdtjd }|d || ||| tj|| W d    n1 s^w   Y  qd S )NrE   void(f2[:], f2)rH   r   rI   )r   r   r   r   rR   iaddisubimulitruedivr_   r`   r   rJ   r=   r>   rc   rb   rL   rM   rd   )	rA   rf   rg   rh   rF   ri   rj   rm   argr   r   r   test_fp16_inplace_binary   s    
z+TestOperatorModule.test_fp16_inplace_binaryc           	   	   C   s   t tf}tjtjf}t||D ]G\}}| j|d5 td|}t	j
dt	jd}t	jdt	j}|d ||d  ||}t	j|| W d    n1 sQw   Y  qd S )NrE   r   rH   r]   rI   r   )r   r   rR   negr   r_   r`   r   rJ   r=   ra   rb   r>   rc   rM   rd   )	rA   rf   rg   rh   rF   ri   rj   rk   rm   r   r   r   test_fp16_unary   s   z"TestOperatorModule.test_fp16_unaryc                 C   2   t d d  t f}tt|dd\}}| d| d S )Nrq   rt   zneg.f16)r   r	   r   rv   rA   ry   rz   r{   r   r   r   test_fp16_neg_ptx   s   z$TestOperatorModule.test_fp16_neg_ptxc                 C   r   )Nrq   rt   zabs.f16)r   r	   r   rv   r   r   r   r   test_fp16_abs_ptx   s   z$TestOperatorModule.test_fp16_abs_ptxc           
   	   C   s   t tttttf}tjtjtj	tj
tjtjf}t||D ]V\}}| j|dD td|}tjdtjd}tjdtj}tjdtj}|d ||d |d  |||}	| |d |	 W d    n1 slw   Y  qd S )NrE   zvoid(b1[:], f2, f2)rH   r]   rI   r   )r    r!   r"   r#   r$   r%   rR   gtgeltleeqner_   r`   r   rJ   r=   ra   bool_r>   rc   rb   assertEqualre   r   r   r   test_fp16_comparison   s&   
z'TestOperatorModule.test_fp16_comparisonc              	   C   s  t tttttf}tjtjtj	tj
tjtjf}tjtjtjtjtjtjf}tt|||D ]X\\}}}| j||dC t|}tjdtjd}tjdtj}	tjdd |}
|d ||	d |
d  ||	|
}| |d | W d    n1 sw   Y  q-d S )Nr}   rH   r]   r   rI   r   ) r    r!   r"   r#   r$   r%   rR   r   r   r   r   r   r   r=   r   r   r   r   r   r   r   r   r_   r`   r   rJ   ra   r   r>   rc   rb   r   )rA   rf   rg   r   rh   rF   r~   ri   rj   rk   rl   rm   r   r   r   test_mixed_fp16_comparison  s0   

z-TestOperatorModule.test_mixed_fp16_comparisonc              	   C      t ttttf}|D ]F}| j|d6 td|}tj	dtj
d}td}td}td}|d |||| | |d	  W d    n1 sJw   Y  q	d S )
Nrh   void(b1[:], f2, f2, f2)rH   r]          @      @g      @rI   r   )r2   r4   r6   r7   r9   r`   r   rJ   r=   ra   r   rb   
assertTruerA   rf   rh   compiledr   rk   rl   arg3r   r   r   !test_multiple_float16_comparisons'  $   


z4TestOperatorModule.test_multiple_float16_comparisonsc              	   C   r   )
Nr   r   rH   r]   r   r   g      ?rI   r   )r2   r4   r6   r7   r9   r`   r   rJ   r=   ra   r   rb   assertFalser   r   r   r   'test_multiple_float16_comparisons_false8  r   z:TestOperatorModule.test_multiple_float16_comparisons_falsec           
   	   C   s   t tttttf}tjtjtj	tj
tjtjf}d}td d  ttf}t|||D ]*\}}}| j|d t||dd\}}	| || W d    n1 sLw   Y  q'd S )N)setp.gt.f16setp.ge.f16setp.lt.f16setp.le.f16setp.eq.f16setp.ne.f16rE   rq   rt   )r    r!   r"   r#   r$   r%   rR   r   r   r   r   r   r   r   r   r_   r`   r	   rv   )
rA   rf   rg   opstringry   rh   rF   srz   r{   r   r   r   test_fp16_comparison_ptxI  s    z+TestOperatorModule.test_fp16_comparison_ptxc           	      C   s   t tttttf}tjtjtj	tj
tjtjf}tjdtjdtj	dtj
dtjdtjdi}t||D ]7\}}| j|d% td d  tttjf}t||dd	\}}| || | W d    n1 saw   Y  q/d S )
Nr   r   r   r   r   r   rE   rq   rt   )r    r!   r"   r#   r$   r%   rR   r   r   r   r   r   r   r_   r`   r   r   r
   r=   r   r	   rv   )	rA   rf   rg   r   rh   rF   ry   rz   r{   r   r   r   test_fp16_int8_comparison_ptxZ  s,   z0TestOperatorModule.test_fp16_int8_comparison_ptxc                 C   s<  t tttttf}tjtjtj	tj
tjtjf}tjtjtjtjtjf}tjdtjdtj	dtj
dtjdtjdi}tddtd	dtd
dtddi}tt|||D ]F\\}}}| j||d1 ttj|}	td d  tt|	f}
t||
dd\}}|| ||	  }| || W d    n1 sw   Y  qUd S )Nzsetp.gt.zsetp.ge.zsetp.lt.zsetp.le.zsetp.eq.z	setp.neu.r   f64r   r   f32r   r}   rq   rt   )r    r!   r"   r#   r$   r%   rR   r   r   r   r   r   r   r=   r   r   r   r   r   r^   r   r   r_   r`   r   rb   r   r   r
   r	   rv   )rA   rf   rg   types_promoter   opsuffixrh   rF   r~   arg2_tyry   rz   r{   r   r   r   (test_mixed_fp16_comparison_promotion_ptxp  sD   



z;TestOperatorModule.test_mixed_fp16_comparison_promotion_ptx)__name__
__module____qualname__r<   rP   rT   rV   rX   rZ   r\   r   rn   r   r|   r   r   r   r   r   r   r   r   r   r   r   r   r   __classcell__r   r   rB   r   r:   o   sN    

	












r:   __main__)-numpyr=   numba.cuda.testingr   r   r   r   numbar   numba.core.typesr   r   
numba.cudar	   rR   r   numba.np.numpy_supportr
   r   r   r   r   r   r   r   r   r   r   r    r!   r"   r#   r$   r%   rJ   r,   r.   r2   r4   r6   r7   r9   r:   r   mainr   r   r   r   <module>   sN    

  #