U
    ?h3                     @   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rZe,  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   W/var/www/html/venv/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_operator.pysimple_fp16_div_scalar   s    r   c                 C   s   || | d< d S r   r   r   r   r   r   simple_fp16add   s    r   c                 C   s   | d  |7  < d S r   r   r   r   r   r   r   simple_fp16_iadd   s    r   c                 C   s   | d  |8  < d S r   r   r   r   r   r   simple_fp16_isub   s    r   c                 C   s   | d  |9  < d S r   r   r   r   r   r   simple_fp16_imul   s    r   c                 C   s   | d  |  < d S r   r   r   r   r   r   simple_fp16_idiv    s    r   c                 C   s   || | d< d S r   r   r   r   r   r   simple_fp16sub$   s    r   c                 C   s   || | d< d S r   r   r   r   r   r   simple_fp16mul(   s    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   s    r   c                 C   s   ||k| d< d S r   r   r   r   r   r   simple_fp16_gt4   s    r   c                 C   s   ||k| d< d S r   r   r   r   r   r   simple_fp16_ge8   s    r   c                 C   s   ||k | d< d S r   r   r   r   r   r   simple_fp16_lt<   s    r    c                 C   s   ||k| d< d S r   r   r   r   r   r   simple_fp16_le@   s    r!   c                 C   s   ||k| d< d S r   r   r   r   r   r   simple_fp16_eqD   s    r"   c                 C   s   ||k| d< d S r   r   r   r   r   r   simple_fp16_neH   s    r#   z
b1(f2, f2)T)Zdevicec                 C   s   | |k S Nr   xyr   r   r   
hlt_func_1L   s    r(   c                 C   s   | |k S r$   r   r%   r   r   r   
hlt_func_2Q   s    r)   c                 C   s   t ||ot||| d< d S r   )r(   r)   rr   r   cr   r   r   test_multiple_hcmp_1V   s    r-   c                 C   s   t ||o||k | d< d S r   r(   r*   r   r   r   test_multiple_hcmp_2[   s    r/   c                 C   s   t ||o||k| d< d S r   r.   r*   r   r   r   test_multiple_hcmp_3`   s    r0   c                 C   s   ||k o||k | d< d S r   r   r*   r   r   r   test_multiple_hcmp_4e   s    r1   c                 C   s   ||k o||k| d< d S r   r   r*   r   r   r   test_multiple_hcmp_5j   s    r2   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   r5   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   rA   rA   )r   jitr6   ZonescopytestingZassert_equal)r:   r?   r@   r   r   resr   r>   r   operator_templatew   s    

z$TestOperatorModule.operator_templatec                 C   s   |  tj d S r$   )rG   operatoraddr9   r   r   r   test_add   s    zTestOperatorModule.test_addc                 C   s   |  tj d S r$   )rG   rH   subr9   r   r   r   test_sub   s    zTestOperatorModule.test_subc                 C   s   |  tj d S r$   )rG   rH   mulr9   r   r   r   test_mul   s    zTestOperatorModule.test_mulc                 C   s   |  tj d S r$   )rG   rH   truedivr9   r   r   r   test_truediv   s    zTestOperatorModule.test_truedivc                 C   s   |  tj d S r$   )rG   rH   floordivr9   r   r   r   test_floordiv   s    z TestOperatorModule.test_floordivc           
   
   C   s   t tttf}tjtjtjtjf}t	||D ]\}}| j
|d| td|}tjdtjd}tjdtj}tjdtj}|d ||d |d  |||}	tj||	 W 5 Q R X q*d S )Nr>   zvoid(f2[:], f2, f2)rA   dtyperB   r   )r   r   r   r   rH   rI   rK   rM   rO   zipsubTestr   rC   r6   zerosfloat16r7   astyperE   assert_allclose
r:   	functionsopsfnr?   kernelgotarg1arg2expectedr   r   r   test_fp16_binary   s    
z#TestOperatorModule.test_fp16_binaryz(Compilation unsupported in the simulatorc              
   C   sn   t ttf}d}td d  ttf}t||D ]>\}}| j|d$ t||dd\}}| || W 5 Q R X q*d S N)zadd.f16zsub.f16zmul.f16)instr      cc)r   r   r   r   rU   rV   r	   assertInr:   r\   Zinstrsargsr^   rf   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 t|}t	jdt	j}t	jdd |}	t	t	j|}
t	jd|
d}|d ||d |	d  |||	}t	j|| W 5 Q R X qNd S )Nr?   tyrA   d   rS   rB   r   )r   r   r   r   rH   rI   rK   rM   rO   r6   int8int16int32int64float32float64	itertoolsproductrU   rV   r   rC   r7   rY   rX   result_typerW   rE   rZ   )r:   r\   r]   typesr^   r?   rs   r_   ra   rb   Zres_tyr`   rc   r   r   r   !test_mixed_fp16_binary_arithmetic   s$     

z4TestOperatorModule.test_mixed_fp16_binary_arithmeticc              
   C   sl   t ttf}d}td d  tf}t||D ]>\}}| j|d$ t||dd\}}| || W 5 Q R X q(d S re   )r   r   r   r   rU   rV   r	   rl   rm   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 ]\}}| j
|dn td|}tjdtj}| }tjdtjd }|d || ||| tj|| W 5 Q R X q*d S )Nr>   void(f2[:], f2)rA   r   rB   )r   r   r   r   rH   iaddisubimulitruedivrU   rV   r   rC   r6   r7   rY   rX   rD   rE   rZ   )	r:   r\   r]   r^   r?   r_   r`   rc   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 ]z\}}| j|d` td|}t	j
dt	jd}t	jdt	j}|d ||d  ||}t	j|| W 5 Q R X qd S )Nr>   r   rA   rS   rB   r   )r   r   rH   negr   rU   rV   r   rC   r6   rW   rX   r7   rY   rE   rZ   )	r:   r\   r]   r^   r?   r_   r`   ra   rc   r   r   r   test_fp16_unary   s    z"TestOperatorModule.test_fp16_unaryc                 C   s2   t d d  t f}tt|dd\}}| d| d S )Nrg   rj   zneg.f16)r   r	   r   rl   r:   rn   ro   rp   r   r   r   test_fp16_neg_ptx   s    z$TestOperatorModule.test_fp16_neg_ptxc                 C   s2   t d d  t f}tt|dd\}}| d| d S )Nrg   rj   zabs.f16)r   r	   r   rl   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 ]\}}| j|d~ td|}tjdtjd}tjdtj}tjdtj}|d ||d |d  |||}	| |d |	 W 5 Q R X q6d S )Nr>   zvoid(b1[:], f2, f2)rA   rS   rB   r   )r   r   r    r!   r"   r#   rH   gtgeltleeqnerU   rV   r   rC   r6   rW   bool8r7   rY   rX   assertEqualr[   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 ]\\}}}| j||d| t|}tjdtjd}tjdtj}	tjdd |}
|d ||	d |
d  ||	|
}| |d | W 5 Q R X qZd S )Nrr   rA   rS   rt   rB   r   ) r   r   r    r!   r"   r#   rH   r   r   r   r   r   r   r6   ru   rv   rw   rx   ry   rz   r{   r|   rU   rV   r   rC   rW   r   r7   rY   rX   r   )r:   r\   r]   r~   r^   r?   rs   r_   r`   ra   rb   rc   r   r   r   test_mixed_fp16_comparison  s2        

z-TestOperatorModule.test_mixed_fp16_comparisonc              
   C   s   t ttttf}|D ]x}| j|db td|}tj	dtj
d}td}td}td}|d |||| | |d	  W 5 Q R X qd S )
Nr^   void(b1[:], f2, f2, f2)rA   rS          @      @g      @rB   r   )r-   r/   r0   r1   r2   rV   r   rC   r6   rW   r   rX   Z
assertTruer:   r\   r^   compiledr   ra   rb   Zarg3r   r   r   !test_multiple_float16_comparisons'  s    


z4TestOperatorModule.test_multiple_float16_comparisonsc              
   C   s   t ttttf}|D ]x}| j|db td|}tj	dtj
d}td}td}td}|d |||| | |d	  W 5 Q R X qd S )
Nr   r   rA   rS   r   r   g      ?rB   r   )r-   r/   r0   r1   r2   rV   r   rC   r6   rW   r   rX   ZassertFalser   r   r   r   'test_multiple_float16_comparisons_false8  s    


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 5 Q R X qNd S )N)setp.gt.f16setp.ge.f16setp.lt.f16setp.le.f16setp.eq.f16setp.ne.f16r>   rg   rj   )r   r   r    r!   r"   r#   rH   r   r   r   r   r   r   r   r   rU   rV   r	   rl   )
r:   r\   r]   opstringrn   r^   r?   sro   rp   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 ]Z\}}| j|d@ td d  tttjf}t||dd	\}}| || | W 5 Q R X q^d S )
Nr   r   r   r   r   r   r>   rg   rj   )r   r   r    r!   r"   r#   rH   r   r   r   r   r   r   rU   rV   r   r   r
   r6   ru   r	   rl   )	r:   r\   r]   r   r^   r?   rn   ro   rp   r   r   r   test_fp16_int8_comparison_ptxZ  s6            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 ]x\\}}}| j||dX ttj|}	td d  tt|	f}
t||
dd\}}|| ||	  }| || W 5 Q R X qd S )Nzsetp.gt.zsetp.ge.zsetp.lt.zsetp.le.zsetp.eq.z	setp.neu.rw   Zf64rx   ry   Zf32rz   rr   rg   rj   )r   r   r    r!   r"   r#   rH   r   r   r   r   r   r   r6   rv   rw   rx   ry   rz   rT   r{   r|   rU   rV   r}   rX   r   r   r
   r	   rl   )r:   r\   r]   Ztypes_promoter   Zopsuffixr^   r?   rs   Zarg2_tyrn   ro   rp   r   r   r   (test_mixed_fp16_comparison_promotion_ptxp  sV             
   z;TestOperatorModule.test_mixed_fp16_comparison_promotion_ptx)__name__
__module____qualname__r5   rG   rJ   rL   rN   rP   rR   r   rd   r   rq   r   r   r   r   r   r   r   r   r   r   r   r   r   __classcell__r   r   r;   r   r3   o   sJ   

	












r3   __main__)-numpyr6   Znumba.cuda.testingr   r   r   r   Znumbar   Znumba.core.typesr   r   Z
numba.cudar	   rH   r{   Znumba.np.numpy_supportr
   r   r   r   r   r   r   r   r   r   r   r   r   r    r!   r"   r#   rC   r(   r)   r-   r/   r0   r1   r2   r3   r   mainr   r   r   r   <module>   sJ   

  #
