U
    b»{jê…  ã                   @   s¾  d dl Z d dlZd dlZd dlZd dlmZm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„ 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+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@„ Z0dAdB„ Z1ej2dCdD�dEdF„ ƒZ3ej2dCdD�dGdH„ ƒZ4dIdJ„ Z5dKdL„ Z6dMdN„ Z7dOdP„ Z8dQdR„ Z9dSdT„ Z:dUdV„ Z;dWdX„ Z<dYdZ„ Z=d[d\„ Z>d]d^„ Z?d_d`„ Z@dadb„ ZAdcdd„ ZBdedf„ ZCdgdh„ ZDdidj„ ZEdkdl„ ZFdmdn„ ZGdodp„ ZHdqdr„ ZIdsdt„ ZJdudv„ ZKdwdx„ ZLdydz„ ZMd{d|„ ZNd}d~„ ZOdd€„ ZPd�d‚„ ZQdƒd„„ ZRd…d†„ ZSd‡dˆ„ ZTd‰dŠ„ ZUd‹dŒ„ ZVG d�dŽ„ dŽeƒZWeXd�k�rºe Y¡  dS )�é    N)ÚcudaÚint64)Úcompile_ptx)ÚTypingError)Úf2)ÚunittestÚCUDATestCaseÚskip_on_cudasimÚskip_unless_cc_53Úskip_unless_cuda_pythonc                 C   s   t jj}|| d< d S ©Nr   ©r   Ú	threadIdxÚx©ÚaryÚi© r   úf/var/www/html/bwave/python/venv/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyÚsimple_threadidx   s    r   c                 C   s   t jj}|| |< d S ©Nr   r   r   r   r   Úfill_threadidx   s    r   c                 C   s>   t jj}t jj}t jj}|d |d  |d  | |||f< d S ©Né   )r   r   r   ÚyÚz)r   r   ÚjÚkr   r   r   Úfill3d_threadidx   s    r   c                 C   s   t  d¡}|| |< d S r   ©r   Úgridr   r   r   r   Úsimple_grid1d   s    
r!   c                 C   s"   t  d¡\}}|| | ||f< d S ©Né   r   )r   r   r   r   r   r   Úsimple_grid2d$   s    r$   c                 C   s(   t  d¡}t  d¡}|dkr$|| d< d S )Nr   r   ©r   r    Úgridsize)r   r   r   r   r   r   Úsimple_gridsize1d)   s    

r'   c                 C   s@   t  d¡\}}t  d¡\}}|dkr<|dkr<|| d< || d< d S )Nr#   r   r   r%   )r   r   r   r   r   r   r   r   Úsimple_gridsize2d0   s
    r(   c           	      C   sp   t  d¡\}}t jjt jj }t jjt jj }| j\}}t|||ƒD ]&}t|||ƒD ]}|| | ||f< qTqDd S r"   )r   r    ÚgridDimr   ÚblockDimr   ÚshapeÚrange)	ÚcÚstartXÚstartYÚgridXÚgridYÚheightÚwidthr   r   r   r   r   Úintrinsic_forloop_step8   s    
r4   c                 C   s   t  |¡| d< d S r   )r   Zpopc©r   r-   r   r   r   Úsimple_popcC   s    r6   c                 C   s   t  |||¡| d< d S r   )r   Úfma©r   ÚaÚbr-   r   r   r   Ú
simple_fmaG   s    r;   c                 C   s   t j |d |d ¡| d< d S r   ©r   Úfp16Zhadd©r   r9   r:   r   r   r   Úsimple_haddK   s    r?   c                 C   s   t j ||¡| d< d S r   r<   r>   r   r   r   Úsimple_hadd_scalarO   s    r@   c                 C   s$   t j |d |d |d ¡| d< d S r   ©r   r=   Zhfmar8   r   r   r   Úsimple_hfmaS   s    rB   c                 C   s   t j |||¡| d< d S r   rA   r8   r   r   r   Úsimple_hfma_scalarW   s    rC   c                 C   s   t j |d |d ¡| d< d S r   ©r   r=   Zhsubr>   r   r   r   Úsimple_hsub[   s    rE   c                 C   s   t j ||¡| d< d S r   rD   r>   r   r   r   Úsimple_hsub_scalar_   s    rF   c                 C   s   t j |d |d ¡| d< d S r   ©r   r=   Zhmulr>   r   r   r   Úsimple_hmulc   s    rH   c                 C   s   t j ||¡| d< d S r   rG   r>   r   r   r   Úsimple_hmul_scalarg   s    rI   c                 C   s   t j ||¡| d< d S r   )r   r=   Úhdivr>   r   r   r   Úsimple_hdiv_scalark   s    rK   c                 C   s:   t  d¡}|| jk r6|| }|| }t j ||¡| |< d S r   )r   r    Úsizer=   rJ   )r   Zarray_aZarray_br   r9   r:   r   r   r   Úsimple_hdiv_kernelo   s
    

rM   c                 C   s   t j |d ¡| d< d S r   ©r   r=   Zhneg©r   r9   r   r   r   Úsimple_hnegw   s    rP   c                 C   s   t j |¡| d< d S r   rN   rO   r   r   r   Úsimple_hneg_scalar{   s    rQ   c                 C   s   t j |d ¡| d< d S r   ©r   r=   ZhabsrO   r   r   r   Úsimple_habs   s    rS   c                 C   s   t j |¡| d< d S r   rR   rO   r   r   r   Úsimple_habs_scalarƒ   s    rT   c                 C   s   t j ||¡| d< d S r   )r   r=   Zheqr>   r   r   r   Úsimple_heq_scalar‡   s    rU   c                 C   s   t j ||¡| d< d S r   )r   r=   Úhner>   r   r   r   Úsimple_hne_scalar‹   s    rW   c                 C   s   t j ||¡| d< d S r   )r   r=   Úhger>   r   r   r   Úsimple_hge_scalar�   s    rY   c                 C   s   t j ||¡| d< d S r   )r   r=   Zhgtr>   r   r   r   Úsimple_hgt_scalar“   s    rZ   c                 C   s   t j ||¡| d< d S r   )r   r=   Zhler>   r   r   r   Úsimple_hle_scalar—   s    r[   c                 C   s   t j ||¡| d< d S r   ©r   r=   Úhltr>   r   r   r   Úsimple_hlt_scalar›   s    r^   T)Zdevicec                 C   s   t j | |¡S r   r\   ©r   r   r   r   r   Ú
hlt_func_1Ÿ   s    r`   c                 C   s   t j | |¡S r   r\   r_   r   r   r   Ú
hlt_func_2¤   s    ra   c                 C   s   t ||ƒot||ƒ| d< d S r   )r`   ra   ©Úrr9   r:   r-   r   r   r   Útest_multiple_hcmp_1©   s    rd   c                 C   s    t ||ƒotj ||¡| d< d S r   )r`   r   r=   r]   rb   r   r   r   Útest_multiple_hcmp_2®   s    re   c                 C   s    t ||ƒotj ||¡| d< d S r   )r`   r   r=   rX   rb   r   r   r   Útest_multiple_hcmp_3³   s    rf   c                 C   s$   t j ||¡ot j ||¡| d< d S r   r\   rb   r   r   r   Útest_multiple_hcmp_4¸   s    rg   c                 C   s$   t j ||¡ot j ||¡| d< d S r   )r   r=   r]   rX   rb   r   r   r   Útest_multiple_hcmp_5½   s    rh   c                 C   s   t j ||¡| d< d S r   )r   r=   Zhmaxr>   r   r   r   Úsimple_hmax_scalarÂ   s    ri   c                 C   s   t j ||¡| d< d S r   )r   r=   Zhminr>   r   r   r   Úsimple_hmin_scalarÆ   s    rj   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    Úlenr=   Zhsin©rc   r   r   r   r   r   Úsimple_hsinÊ   s    
rm   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhcosrl   r   r   r   Úsimple_hcosÑ   s    
rn   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhlogrl   r   r   r   Úsimple_hlogØ   s    
ro   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhlog2rl   r   r   r   Úsimple_hlog2ß   s    
rp   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhlog10rl   r   r   r   Úsimple_hlog10æ   s    
rq   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhexprl   r   r   r   Úsimple_hexpí   s    
rr   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhexp2rl   r   r   r   Úsimple_hexp2ô   s    
rs   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhsqrtrl   r   r   r   Úsimple_hsqrtû   s    
rt   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhrsqrtrl   r   r   r   Úsimple_hrsqrt  s    
ru   c                 C   s   | d S )Ng      à¿r   )r   Údtyper   r   r   Únumpy_hrsqrt
  s    rw   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhceilrl   r   r   r   Úsimple_hceil  s    
rx   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhfloorrl   r   r   r   Úsimple_hfloor  s    
ry   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhrcprl   r   r   r   Úsimple_hrcp  s    
rz   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhtruncrl   r   r   r   Úsimple_htrunc#  s    
r{   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhrintrl   r   r   r   Úsimple_hrint*  s    
r|   c                 C   s   t  |¡| d< d S r   )r   ZcbrtrO   r   r   r   Úsimple_cbrt1  s    r}   c                 C   s   t  |¡| d< d S r   )r   Zbrevr5   r   r   r   Úsimple_brev5  s    r~   c                 C   s   t  |¡| d< d S r   )r   Zclzr5   r   r   r   Ú
simple_clz9  s    r   c                 C   s   t  |¡| d< d S r   )r   Zffsr5   r   r   r   Ú
simple_ffs=  s    r€   c                 C   s   t |ƒ| d< d S r   ©Úroundr5   r   r   r   Úsimple_roundA  s    rƒ   c                 C   s   t ||ƒ| d< d S r   r�   )r   r-   Úndigitsr   r   r   Úsimple_round_toE  s    r…   c                 C   sF   t  d¡}| | dkr:|d dkr0|| | |< qBd| |< nd| |< d S )Nr   é   r#   r   é   é   r   )r9   r:   r-   r   r   r   r   Úbranching_with_ifsI  s    

r‰   c                 C   sB   t  d¡}t  |d dk|| d¡}t  | | dk|d¡| |< d S )Nr   r#   r   r‡   r†   rˆ   )r   r    Úselp)r9   r:   r-   r   Úinnerr   r   r   Úbranching_with_selpsU  s    
rŒ   c                 C   s   t  d¡}t j| |< d S r   )r   r    Zlaneidr   r   r   r   Úsimple_laneid\  s    
r�   c                 C   s   t j| d< d S r   )r   Zwarpsize)r   r   r   r   Úsimple_warpsizea  s    rŽ   c                 C   s   t  | ¡ d S r   r   ©r   r   r   r   Únonliteral_gride  s    r�   c                 C   s   t  | ¡ d S r   )r   r&   r�   r   r   r   Únonliteral_gridsizei  s    r‘   c                       s  e Zd Z‡ fdd„Zdd„ Zdd„ Zdd„ Zed	ƒd
d„ ƒZed	ƒdd„ ƒZ	dd„ Z
dd„ Zdd„ Ze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ed'd(„ ƒZed)d*„ ƒZed+ƒd,d-„ ƒZed.d/„ ƒZed0d1„ ƒZed+ƒd2d3„ ƒZed4d5„ ƒZed6d7„ ƒZed+ƒd8d9„ ƒZed:d;„ ƒZ ed<d=„ ƒZ!ed+ƒd>d?„ ƒZ"ee#d@ƒdAdB„ ƒƒZ$ee#d@ƒdCdD„ ƒƒZ%edEdF„ ƒZ&edGdH„ ƒZ'ed+ƒdIdJ„ ƒZ(edKdL„ ƒZ)edMdN„ ƒZ*ed+ƒdOdP„ ƒZ+ee#d@ƒdQdR„ ƒƒZ,ee#d@ƒdSdT„ ƒƒZ-edUdV„ ƒZ.edWdX„ ƒZ/edYdZ„ ƒZ0ed[d\„ ƒZ1d]d^„ Z2d_d`„ Z3dadb„ Z4edcƒddde„ ƒZ5dfdg„ Z6dhdi„ Z7djdk„ Z8dldm„ Z9edcƒdndo„ ƒZ:dpdq„ Z;drds„ Z<dtdu„ Z=dvdw„ Z>edcƒdxdy„ ƒZ?dzd{„ Z@d|d}„ ZAd~d„ ZBd€d�„ ZCd‚dƒ„ ZDed„ƒd…d†„ ƒZEd‡dˆ„ ZFd‰dŠ„ ZGed„ƒd‹dŒ„ ƒZHd�dŽ„ ZI‡  ZJS )�ÚTestCudaIntrinsicc                    s   t ƒ  ¡  tj d¡ d S r   )ÚsuperÚsetUpÚnpÚrandomÚseed©Úself©Ú	__class__r   r   r”   n  s    
zTestCudaIntrinsic.setUpc                 C   s@   t  d¡tƒ}tjdtjd�}|d |ƒ |  |d dk¡ d S )Núvoid(int32[:])r   ©rv   ©r   r   r   )r   Újitr   r•   ÚonesÚint32Ú
assertTrue©r™   Úcompiledr   r   r   r   Útest_simple_threadidxr  s    z'TestCudaIntrinsic.test_simple_threadidxc                 C   sZ   t  d¡tƒ}d}tj|tjd�}tj|tjd�}|d|f |ƒ |  t ||k¡¡ d S )Nrœ   é
   r�   r   )	r   rŸ   r   r•   r    r¡   Úaranger¢   Úall)r™   r¤   ÚNr   Úexpr   r   r   Útest_fill_threadidxx  s    z%TestCudaIntrinsic.test_fill_threadidxc                    sN   d\‰ ‰‰‡ ‡‡fdd„}‡ ‡‡fdd„}|ƒ }|ƒ }|   t ||k¡¡ d S )N)r†   é   é   c                     s>   t  d¡tƒ} tjˆ ˆˆftjd�}| dˆ ˆˆff |ƒ |S )Nzvoid(int32[:,:,::1])r�   r   )r   rŸ   r   r•   Úzerosr¡   ©r¤   r   ©ÚXÚYÚZr   r   Úc_contigousƒ  s    z<TestCudaIntrinsic.test_fill3d_threadidx.<locals>.c_contigousc                     sD   t  d¡tƒ} t tjˆ ˆˆftjd�¡}| dˆ ˆˆff |ƒ |S )Nzvoid(int32[::1,:,:])r�   r   )r   rŸ   r   r•   Zasfortranarrayr®   r¡   r¯   r°   r   r   Úf_contigous‰  s    z<TestCudaIntrinsic.test_fill3d_threadidx.<locals>.f_contigous)r¢   r•   r¨   )r™   r´   rµ   Zc_resZf_resr   r°   r   Útest_fill3d_threadidx€  s    
z'TestCudaIntrinsic.test_fill3d_threadidxzCudasim does not check typesc              	   C   s*   |   td¡� t d¡tƒ W 5 Q R X d S ©NZRequireLiteralValuezvoid(int32))ÚassertRaisesRegexr   r   rŸ   r�   r˜   r   r   r   Útest_nonliteral_grid_error“  s    z,TestCudaIntrinsic.test_nonliteral_grid_errorc              	   C   s*   |   td¡� t d¡tƒ W 5 Q R X d S r·   )r¸   r   r   rŸ   r‘   r˜   r   r   r   Útest_nonliteral_gridsize_error˜  s    z0TestCudaIntrinsic.test_nonliteral_gridsize_errorc                 C   s\   t  d¡tƒ}d\}}|| }tj|tjd�}|||f |ƒ |  t |t |¡k¡¡ d S )Núvoid(int32[::1])©rˆ   é   r�   )	r   rŸ   r!   r•   Úemptyr¡   r¢   r¨   r§   )r™   r¤   ÚntidÚnctaidZnelemr   r   r   r   Útest_simple_grid1d�  s    z$TestCudaIntrinsic.test_simple_grid1dc           	      C   s®   t  d¡tƒ}d}d}|d |d  |d |d  f}tj|tjd�}| ¡ }|||f |ƒ t|jd ƒD ](}t|jd ƒD ]}|| |||f< q~ql|  	t 
||k¡¡ d S ©Nzvoid(int32[:,::1])©r†   rˆ   ©r¬   r­   r   r   r�   )r   rŸ   r$   r•   r¾   r¡   Úcopyr,   r+   r¢   r¨   )	r™   r¤   r¿   rÀ   r+   r   rª   r   r   r   r   r   Útest_simple_grid2d¥  s     z$TestCudaIntrinsic.test_simple_grid2dc                 C   sN   t  d¡tƒ}d\}}tjdtjd�}|||f |ƒ |  |d || ¡ d S )Nr»   r¼   r   r�   r   )r   rŸ   r'   r•   r®   r¡   ÚassertEqual©r™   r¤   r¿   rÀ   r   r   r   r   Útest_simple_gridsize1d´  s
    z(TestCudaIntrinsic.test_simple_gridsize1dzTests PTX emissionc           
      C   s  t d d … t t d d … f}t |¡tƒ}t |¡tƒ}d}d}tjddtj d�}| ¡ }d|d d…< tj|tj d�}||df |||ƒ | 	|¡}	|  
d	tt d
|	¡ƒ¡ tjj||dd� tj|tj d�}||df |||ƒ | 	|¡}	|  
dtt d
|	¡ƒ¡ tjj||dd� d S )Né    r­   é   )r+   Z
fill_valuerv   rˆ   r¬   r�   r   r#   z	\s+bra\s+Z	branching)Úerr_msgr   rŠ   )r   r   rŸ   r‰   rŒ   r•   ÚfullrÅ   r§   Zinspect_asmrÇ   rk   ÚreÚfindallÚtestingZassert_array_equal)
r™   ÚsigZcu_branching_with_ifsZcu_branching_with_selpsÚnr:   r-   Úexpectedr9   Úptxr   r   r   Ú	test_selp»  s$    

zTestCudaIntrinsic.test_selpc                 C   sr   t  d¡tƒ}d}d}tjdtjd�}|||f |ƒ |  |d |d |d  ¡ |  |d |d |d  ¡ d S )Nr»   rÃ   rÄ   r#   r�   r   r   )r   rŸ   r(   r•   r®   r¡   rÇ   rÈ   r   r   r   Útest_simple_gridsize2dÔ  s    z(TestCudaIntrinsic.test_simple_gridsize2dc              	   C   sò   t  d¡tƒ}d}d}|d |d  |d |d  f}tj|tjd�}|||f |ƒ |\}}|j\}}	tt|d ƒt|d ƒƒD ]j\}
}||
 ||  }}t||	|ƒD ]B}t|||ƒD ]0}|  	|||f || k|||f || f¡ q¸q¨q‚d S rÂ   )
r   rŸ   r4   r•   r¾   r¡   r+   Úzipr,   r¢   )r™   r¤   r¿   rÀ   r+   r   r0   r1   r2   r3   r   r   r.   r/   r   r   r   r   r   Útest_intrinsic_forloop_stepÞ  s     
"z-TestCudaIntrinsic.test_intrinsic_forloop_stepc                 C   sF   t jdd„ ƒ}tjdtjd� ddd¡}|d |ƒ tj |d¡ d S )Nc                 S   s:   t  d¡\}}}t  d¡\}}}|| | | |||f< d S ©Nrˆ   r%   )Úoutr   r   r   r9   r:   r-   r   r   r   Úfooð  s    z*TestCudaIntrinsic.test_3dgrid.<locals>.fooiÙ  r�   é	   )©rˆ   rˆ   rˆ   rÝ   )r   rŸ   r•   r®   r¡   ÚreshaperÐ   Zassert_equal)r™   rÛ   Úarrr   r   r   Útest_3dgridï  s
    
zTestCudaIntrinsic.test_3dgridc                 C   sZ   t jdd„ ƒ}d\}}}tj|| | tjd� |||¡}|d |ƒ |  t |¡¡ d S )Nc           	      S   sÆ   t  d¡\}}}t  d¡\}}}|t jjt jjt jj  kor|t jjt jjt jj  kor|t jjt jjt jj  k}|t jjt j	j ko®|t jjt j	j ko®|t jjt j	j k}|o¶|| |||f< d S rÙ   )
r   r    r&   r   r   ZblockIdxr*   r   r   r)   )	rÚ   r   r   r   r9   r:   r-   Zgrid_is_rightZgridsize_is_rightr   r   r   rÛ   ü  s    ÿýÿþz,TestCudaIntrinsic.test_3dgrid_2.<locals>.foo)é   r­   é   r�   ))r†   rˆ   r#   )rˆ   r#   r†   )r   rŸ   r•   r®   Zbool_rÞ   r¢   r¨   )r™   rÛ   r   r   r   rß   r   r   r   Útest_3dgrid_2û  s    

"zTestCudaIntrinsic.test_3dgrid_2c                 C   s@   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ d S )Núvoid(int32[:], uint32)r   r�   rž   éð   r   r†   ©r   rŸ   r6   r•   r®   r¡   ÚassertEqualsr£   r   r   r   Útest_popc_u4  s    zTestCudaIntrinsic.test_popc_u4c                 C   s@   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ d S )Nzvoid(int32[:], uint64)r   r�   rž   l        @ r   r†   ræ   r£   r   r   r   Útest_popc_u8  s    zTestCudaIntrinsic.test_popc_u8c                 C   sF   t  d¡tƒ}tjdtjd�}|d |dddƒ tj |d d	¡ d S )
Nzvoid(f4[:], f4, f4, f4)r   r�   rž   ç       @ç      @ç      @r   r¦   )r   rŸ   r;   r•   r®   Úfloat32rÐ   Úassert_allcloser£   r   r   r   Útest_fma_f4  s    zTestCudaIntrinsic.test_fma_f4c                 C   sF   t  d¡tƒ}tjdtjd�}|d |dddƒ tj |d d	¡ d S )
Nzvoid(f8[:], f8, f8, f8)r   r�   rž   rê   rë   rì   r   r¦   )r   rŸ   r;   r•   r®   Úfloat64rÐ   rî   r£   r   r   r   Útest_fma_f8"  s    zTestCudaIntrinsic.test_fma_f8c                 C   sl   t  d¡tƒ}tjdtjd�}tjdgtjd�}tjdgtjd�}|d |||ƒ tj |d || ¡ d S ©Núvoid(f2[:], f2[:], f2[:])r   r�   rë   rì   rž   r   )	r   rŸ   r?   r•   r®   Úfloat16ÚarrayrÐ   rî   ©r™   r¤   r   Úarg1Úarg2r   r   r   Ú	test_hadd(  s    zTestCudaIntrinsic.test_haddc                 C   s`   t  d¡tƒ}tjdtjd�}t d¡}t d¡}|d |||ƒ || }tj |d |¡ d S )Núvoid(f2[:], f2, f2)r   r�   çJØMû!	@rë   rž   r   )r   rŸ   r@   r•   r®   rô   rÐ   rî   ©r™   r¤   r   r÷   rø   Úrefr   r   r   Útest_hadd_scalar1  s    

z"TestCudaIntrinsic.test_hadd_scalarz(Compilation unsupported in the simulatorc                 C   s4   t d d … t t f}tt|dd�\}}|  d|¡ d S )N©r¬   rˆ   ©Úcczadd.f16)r   r   r@   ÚassertIn©r™   ÚargsrÔ   Ú_r   r   r   Útest_hadd_ptx;  s    zTestCudaIntrinsic.test_hadd_ptxc                 C   s„   t  d¡tƒ}tjdtjd�}tjdgtjd�}tjdgtjd�}tjdgtjd�}|d ||||ƒ tj |d || | ¡ d S )	Nz void(f2[:], f2[:], f2[:], f2[:])r   r�   rê   rë   rì   rž   r   )	r   rŸ   rB   r•   r®   rô   rõ   rÐ   rî   )r™   r¤   r   r÷   rø   Úarg3r   r   r   Ú	test_hfmaA  s    zTestCudaIntrinsic.test_hfmac                 C   sp   t  d¡tƒ}tjdtjd�}t d¡}t d¡}t d¡}|d ||||ƒ || | }tj |d |¡ d S )	Nzvoid(f2[:], f2, f2, f2)r   r�   rê   rë   rì   rž   r   )r   rŸ   rC   r•   r®   rô   rÐ   rî   )r™   r¤   r   r÷   rø   r  rý   r   r   r   Útest_hfma_scalarK  s    


z"TestCudaIntrinsic.test_hfma_scalarc                 C   s6   t d d … t t t f}tt|dd�\}}|  d|¡ d S )Nrÿ   r   z
fma.rn.f16)r   r   rC   r  r  r   r   r   Útest_hfma_ptxV  s    zTestCudaIntrinsic.test_hfma_ptxc                 C   sl   t  d¡tƒ}tjdtjd�}tjdgtjd�}tjdgtjd�}|d |||ƒ tj |d || ¡ d S rò   )	r   rŸ   rE   r•   r®   rô   rõ   rÐ   rî   rö   r   r   r   Ú	test_hsub\  s    zTestCudaIntrinsic.test_hsubc                 C   s`   t  d¡tƒ}tjdtjd�}t d¡}t d¡}|d |||ƒ || }tj |d |¡ d S ©Nrú   r   r�   rû   g…ëQ¸ù?rž   r   )r   rŸ   rF   r•   r®   rô   rÐ   rî   rü   r   r   r   Útest_hsub_scalare  s    

z"TestCudaIntrinsic.test_hsub_scalarc                 C   s4   t d d … t t f}tt|dd�\}}|  d|¡ d S )Nrÿ   r   zsub.f16)r   r   rF   r  r  r   r   r   Útest_hsub_ptxo  s    zTestCudaIntrinsic.test_hsub_ptxc                 C   sj   t  ¡ tƒ}tjdtjd�}tjdgtjd�}tjdgtjd�}|d |||ƒ tj |d || ¡ d S )Nr   r�   rë   rì   rž   r   )	r   rŸ   rH   r•   r®   rô   rõ   rÐ   rî   rö   r   r   r   Ú	test_hmulu  s    zTestCudaIntrinsic.test_hmulc                 C   s`   t  d¡tƒ}tjdtjd�}t d¡}t d¡}|d |||ƒ || }tj |d |¡ d S r  )r   rŸ   rI   r•   r®   rô   rÐ   rî   rü   r   r   r   Útest_hmul_scalar~  s    

z"TestCudaIntrinsic.test_hmul_scalarc                 C   s4   t d d … t t f}tt|dd�\}}|  d|¡ d S )Nrÿ   r   zmul.f16)r   r   rI   r  r  r   r   r   Útest_hmul_ptxˆ  s    zTestCudaIntrinsic.test_hmul_ptxzNVIDIA Binding needed for NVRTCc                 C   s`   t  d¡tƒ}tjdtjd�}t d¡}t d¡}|d |||ƒ || }tj |d |¡ d S r  )r   rŸ   rK   r•   r®   rô   rÐ   rî   rü   r   r   r   Útest_hdiv_scalarŽ  s    

z"TestCudaIntrinsic.test_hdiv_scalarc                 C   s€   t  d¡tƒ}tjjdddd� tj¡}tjjdddd� tj¡}tj|tjd�}| 	|j
¡|||ƒ || }tj ||¡ d S )Nró   i  ÿÿéáÿ  iô  ©rL   r�   )r   rŸ   rM   r•   r–   ÚrandintÚastyperô   Ú
zeros_likeÚforallrL   rÐ   rî   )r™   r¤   Zarry1Zarry2r   rý   r   r   r   Ú	test_hdiv™  s    zTestCudaIntrinsic.test_hdivc                 C   sV   t  d¡tƒ}tjdtjd�}tjdgtjd�}|d ||ƒ tj |d | ¡ d S )Núvoid(f2[:], f2[:])r   r�   rë   rž   r   )	r   rŸ   rP   r•   r®   rô   rõ   rÐ   rî   ©r™   r¤   r   r÷   r   r   r   Ú	test_hneg¥  s
    zTestCudaIntrinsic.test_hnegc                 C   sR   t  d¡tƒ}tjdtjd�}t d¡}|d ||ƒ | }tj |d |¡ d S )Núvoid(f2[:], f2)r   r�   rû   rž   r   )r   rŸ   rQ   r•   r®   rô   rÐ   rî   ©r™   r¤   r   r÷   rý   r   r   r   Útest_hneg_scalar­  s    
z"TestCudaIntrinsic.test_hneg_scalarc                 C   s2   t d d … t f}tt|dd�\}}|  d|¡ d S )Nrÿ   r   zneg.f16)r   r   rQ   r  r  r   r   r   Útest_hneg_ptx¶  s    zTestCudaIntrinsic.test_hneg_ptxc                 C   sV   t  ¡ tƒ}tjdtjd�}tjdgtjd�}|d ||ƒ tj |d t	|ƒ¡ d S )Nr   r�   ç      Àrž   r   )
r   rŸ   rS   r•   r®   rô   rõ   rÐ   rî   Úabsr  r   r   r   Ú	test_habs¼  s
    zTestCudaIntrinsic.test_habsc                 C   sT   t  d¡tƒ}tjdtjd�}t d¡}|d ||ƒ t|ƒ}tj |d |¡ d S )Nr  r   r�   gJØMû!	Àrž   r   )	r   rŸ   rT   r•   r®   rô   r"  rÐ   rî   r  r   r   r   Útest_habs_scalarÄ  s    
z"TestCudaIntrinsic.test_habs_scalarc                 C   s2   t d d … t f}tt|dd�\}}|  d|¡ d S )Nrÿ   r   zabs.f16)r   r   rT   r  r  r   r   r   Útest_habs_ptxÍ  s    zTestCudaIntrinsic.test_habs_ptxc                 C   s†  t ttttttttt	t
tf}ttf}tjtjtjtjtjtjtjtjtjtjtjtf}tjtjf}d}tj d¡ tjjdd|d�  tj!¡}t "|¡}t#||ƒD ]\\}}	| j$|	d��B t% &d¡|ƒ}|d|f ||ƒ |	|tj!d�}
tj' (||
¡ W 5 Q R X q tjjdd|d�  tj!¡}t#||ƒD ]^\}}	| j$|	d��B t% &d¡|ƒ}|d|f ||ƒ |	|tj!d�}
tj' (||
¡ W 5 Q R X �q"d S )	NrÊ   r   r  r  ©Úfnr  r�   r¦   ))rm   rn   ro   rp   rq   rt   rx   ry   rz   r{   r|   ru   rr   rs   r•   ÚsinÚcosÚlogÚlog2Úlog10ÚsqrtÚceilÚfloorZ
reciprocalÚtruncZrintrw   rª   Zexp2r–   r—   r  r  rô   r  r×   ÚsubTestr   rŸ   rÐ   rî   )r™   ZkernelsZexp_kernelsZexpected_functionsZexpected_exp_functionsr©   r   rc   Úkernelr'  rÓ   Zx2r   r   r   Útest_fp16_intrinsics_commonÓ  sV          ü      ü
z-TestCudaIntrinsic.test_fp16_intrinsics_commonc                 C   sf   t  ¡ dd„ ƒ}d}tj d¡ tj |¡ tj¡}t |¡}|d|f ||ƒ tj	 
|d| ¡ d S )Nc                 S   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rk   r=   Zhexp10rl   r   r   r   Úhexp10_vectorsú  s    
z5TestCudaIntrinsic.test_hexp10.<locals>.hexp10_vectorsrÊ   r   r¦   )r   rŸ   r•   r–   r—   Zrandr  rô   r  rÐ   rî   )r™   r4  r©   r   rc   r   r   r   Útest_hexp10÷  s    

zTestCudaIntrinsic.test_hexp10c              
   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d�}t d¡}t d¡}	t d¡}
|d ||	|	ƒ ||	|	ƒ}|  ||d	 ¡ |d ||	|
ƒ ||	|
ƒ}|  ||d	 ¡ |d ||	|ƒ ||	|ƒ}|  ||d	 ¡ W 5 Q R X q6d S )
N)Úopzvoid(b1[:], f2, f2)r   r�   r#   rˆ   r†   rž   r   )rU   rW   rY   rZ   r[   r^   ÚoperatorÚeqÚneÚgeÚgtÚleÚltr×   r1  r   rŸ   r•   r®   Úbool8rô   rÇ   )r™   ZfnsZopsr'  r6  r2  rÓ   Úgotrø   r  Zarg4r   r   r   Útest_fp16_comparison  s6      ÿ  ÿ





z&TestCudaIntrinsic.test_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 )
Nr&  zvoid(b1[:], f2, f2, f2)r   r�   rê   rë   rì   rž   r   )rd   re   rf   rg   rh   r1  r   rŸ   r•   r®   r>  rô   r¢   )r™   Z	functionsr'  r¤   r   r÷   rø   r  r   r   r   Ú!test_multiple_float16_comparisons+  s    ü


z3TestCudaIntrinsic.test_multiple_float16_comparisonsc                 C   s„   t  d¡tƒ}tjdtjd�}t d¡}t d¡}|d |||ƒ tj |d |¡ t d¡}|d |||ƒ tj |d |¡ d S ©	Nrú   r   r�   rë   rì   rž   r   g      @)r   rŸ   ri   r•   r®   rô   rÐ   rî   rö   r   r   r   Ú	test_hmax<  s    


zTestCudaIntrinsic.test_hmaxc                 C   s„   t  d¡tƒ}tjdtjd�}t d¡}t d¡}|d |||ƒ tj |d |¡ t d¡}|d |||ƒ tj |d |¡ d S rB  )r   rŸ   rj   r•   r®   rô   rÐ   rî   rö   r   r   r   Ú	test_hminH  s    


zTestCudaIntrinsic.test_hminc                 C   sJ   t  d¡tƒ}tjdtjd�}d}|d ||ƒ tj |d |d ¡ d S )Nzvoid(float32[:], float32)r   r�   rê   rž   r   çUUUUUUÕ?)r   rŸ   r}   r•   r®   rí   rÐ   rî   ©r™   r¤   r   Zcbrt_argr   r   r   Útest_cbrt_f32T  s
    zTestCudaIntrinsic.test_cbrt_f32c                 C   sJ   t  d¡tƒ}tjdtjd�}d}|d ||ƒ tj |d |d ¡ d S )Nzvoid(float64[:], float64)r   r�   g      @rž   r   rE  )r   rŸ   r}   r•   r®   rð   rÐ   rî   rF  r   r   r   Útest_cbrt_f64[  s
    zTestCudaIntrinsic.test_cbrt_f64c                 C   s@   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ d S )Nzvoid(uint32[:], uint32)r   r�   rž   ið0  r   i  )r   rŸ   r~   r•   r®   Zuint32rç   r£   r   r   r   Útest_brev_u4b  s    zTestCudaIntrinsic.test_brev_u4z.only get given a Python "int", assumes 32 bitsc                 C   s@   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ d S )Nzvoid(uint64[:], uint64)r   r�   rž   l   ð0  ÀC r   l       `x)r   rŸ   r~   r•   r®   Zuint64rç   r£   r   r   r   Útest_brev_u8h  s    zTestCudaIntrinsic.test_brev_u8c                 C   s@   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ d S )Núvoid(int32[:], int32)r   r�   rž   é   r   é   ©r   rŸ   r   r•   r®   r¡   rç   r£   r   r   r   Útest_clz_i4o  s    zTestCudaIntrinsic.test_clz_i4c                 C   s@   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ dS )	aº  
        Although the CUDA Math API
        (http://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__INTRINSIC__INT.html)
        only says int32 & int64 arguments are supported in C code, the LLVM
        IR input supports i8, i16, i32 & i64 (LLVM doesn't have a concept of
        unsigned integers, just unsigned operations on integers).
        http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics
        rä   r   r�   rž   rL  r   rM  NrN  r£   r   r   r   Útest_clz_u4u  s    	zTestCudaIntrinsic.test_clz_u4c                 C   s@   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ d S ©NrK  r   r�   rž   l   ÿÿ r   rN  r£   r   r   r   Útest_clz_i4_1sƒ  s    z TestCudaIntrinsic.test_clz_i4_1sc                 C   sB   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d dd¡ d S )NrK  r   r�   rž   r   rÊ   úCUDA semanticsrN  r£   r   r   r   Útest_clz_i4_0s‰  s    z TestCudaIntrinsic.test_clz_i4_0sc                 C   s@   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ d S )Núvoid(int32[:], int64)r   r�   rž   é   r   é/   rN  r£   r   r   r   Útest_clz_i8�  s    zTestCudaIntrinsic.test_clz_i8c                 C   s^   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ |d |dƒ |  |d d	¡ d S )
NrK  r   r�   rž   rL  r   é   ì        rÊ   ©r   rŸ   r€   r•   r®   r¡   rç   r£   r   r   r   Útest_ffs_i4–  s    zTestCudaIntrinsic.test_ffs_i4c                 C   s^   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ |d |dƒ |  |d d	¡ d S )
Nrä   r   r�   rž   rL  r   rY  rZ  rÊ   r[  r£   r   r   r   Útest_ffs_u4ž  s    zTestCudaIntrinsic.test_ffs_u4c                 C   s@   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ d S rQ  r[  r£   r   r   r   Útest_ffs_i4_1s¦  s    z TestCudaIntrinsic.test_ffs_i4_1sc                 C   s@   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ d S )NrK  r   r�   rž   r   r[  r£   r   r   r   Útest_ffs_i4_0s¬  s    z TestCudaIntrinsic.test_ffs_i4_0sc                 C   s^   t  d¡tƒ}tjdtjd�}|d |dƒ |  |d d¡ |d |dƒ |  |d d	¡ d S )
NrU  r   r�   rž   rV  r   rË   l        é!   r[  r£   r   r   r   Útest_ffs_i8²  s    zTestCudaIntrinsic.test_ffs_i8c                 C   sj   t  d¡tƒ}d}tj|d tjd�}t tjdtjd�|¡}|d|d f |ƒ |  t 	||k¡¡ d S )Nrœ   r#   rÊ   r�   r   )
r   rŸ   r�   r•   r®   r¡   Ztiler§   r¢   r¨   )r™   r¤   Úcountr   rª   r   r   r   Útest_simple_laneid»  s    z$TestCudaIntrinsic.test_simple_laneidc                 C   s@   t  d¡tƒ}tjdtjd�}|d |ƒ |  |d dd¡ d S )Nrœ   r   r�   rž   r   rÊ   rS  )r   rŸ   rŽ   r•   r®   r¡   rç   r£   r   r   r   Útest_simple_warpsizeÃ  s    z&TestCudaIntrinsic.test_simple_warpsizec                 C   sN   t  d¡tƒ}tjdtjd�}dD ]&}|d ||ƒ |  |d t|ƒ¡ q"d S )Nzvoid(int64[:], float32)r   r�   ©r!  g      Àg      Àg      ø¿g      ø?g      @g      @g      @rž   r   ©r   rŸ   rƒ   r•   r®   r   rç   r‚   ©r™   r¤   r   r   r   r   r   Útest_round_f4É  s
    zTestCudaIntrinsic.test_round_f4c                 C   sN   t  d¡tƒ}tjdtjd�}dD ]&}|d ||ƒ |  |d t|ƒ¡ q"d S )Nzvoid(int64[:], float64)r   r�   re  rž   r   rf  rg  r   r   r   Útest_round_f8Ñ  s
    zTestCudaIntrinsic.test_round_f8c              
   C   sÂ   t  d¡tƒ}tjdtjd�}tj d¡ tj d¡ tj¡}t 	|t 
tjtj tjg¡f¡ d}t ||¡D ]L\}}| j||d��0 |d |||ƒ | j|d	 t||ƒd
d� W 5 Q R X qpd S )Nú void(float32[:], float32, int32)r   r�   é{   rÊ   )éûÿÿÿéüÿÿÿéýÿÿÿéþÿÿÿéÿÿÿÿr   r   r#   rˆ   r†   r¬   r‡   ©Úvalr„   rž   r   Úsingle©Úprec)r   rŸ   r…   r•   r®   rí   r–   r—   r  Úconcatenaterõ   ÚinfÚnanÚ	itertoolsÚproductr1  ÚassertPreciseEqualr‚   ©r™   r¤   r   ÚvalsÚdigitsrr  r„   r   r   r   Útest_round_to_f4Ù  s    "	ÿz"TestCudaIntrinsic.test_round_to_f4z$Overflow behavior differs on CPythonc                 C   sT   t  d¡tƒ}tjdtjd�}t tj¡j}d}|d |||ƒ |  |d |¡ d S )Nrj  r   r�   i,  rž   r   )	r   rŸ   r…   r•   r®   rí   ÚfinfoÚmaxrÇ   ©r™   r¤   r   rr  r„   r   r   r   Útest_round_to_f4_overflowó  s    z+TestCudaIntrinsic.test_round_to_f4_overflowc                 C   sT   t  d¡tƒ}tjdtjd�}d}d}|d |||ƒ | j|d t||ƒdd	� d S )
Nrj  r   r�   gìQ¸…ëÕ?rˆ   rž   r   rs  rt  )r   rŸ   r…   r•   r®   rí   r{  r‚   r‚  r   r   r   Útest_round_to_f4_halfway  s    z*TestCudaIntrinsic.test_round_to_f4_halfwayc              
   C   s  t  d¡tƒ}tjdtjd�}tj d¡ tj d¡}t |t 	tj
tj
 tjg¡f¡ d}t ||¡D ]L\}}| j||d��0 |d |||ƒ | j|d	 t||ƒd
d� W 5 Q R X qhd}d}| j||d��0 |d |||ƒ | j|d	 t||ƒdd� W 5 Q R X d S )Nú void(float64[:], float64, int32)r   r�   rk  rÊ   )rl  rm  rn  ro  rp  r   r   r#   rˆ   r†   r¬   rq  rž   r   Úexactrt  g`¢8„p=Ö<é   Údouble)r   rŸ   r…   r•   r®   rð   r–   r—   rv  rõ   rw  rx  ry  rz  r1  r{  r‚   r|  r   r   r   Útest_round_to_f8  s&    "ÿÿz"TestCudaIntrinsic.test_round_to_f8c                 C   sT   t  d¡tƒ}tjdtjd�}t tj¡j}d}|d |||ƒ |  |d |¡ d S )Nr…  r   r�   rá   rž   r   )	r   rŸ   r…   r•   r®   rð   r€  r�  rÇ   r‚  r   r   r   Útest_round_to_f8_overflow"  s    z+TestCudaIntrinsic.test_round_to_f8_overflowc                 C   sT   t  d¡tƒ}tjdtjd�}d}d}|d |||ƒ | j|d t||ƒdd	� d S )
Nr…  r   r�   g\�Âõ(\á?rˆ   rž   r   rˆ  rt  )r   rŸ   r…   r•   r®   rð   r{  r‚   r‚  r   r   r   Útest_round_to_f8_halfway/  s    z*TestCudaIntrinsic.test_round_to_f8_halfway)KÚ__name__Ú
__module__Ú__qualname__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  r  r  r   r#  r$  r%  r3  r5  r@  rA  rC  rD  rG  rH  rI  rJ  rO  rP  rR  rT  rX  r\  r]  r^  r_  ra  rc  rd  rh  ri  r  rƒ  r„  r‰  rŠ  r‹  Ú__classcell__r   r   rš   r   r’   m  sÔ   





	

	




	


	
	






"









r’   Ú__main__)Zry  Únumpyr•   r7  rÎ   Znumbar   r   Z
numba.cudar   Znumba.core.errorsr   Znumba.core.typesr   Znumba.cuda.testingr   r   r	   r
   r   r   r   r   r!   r$   r'   r(   r4   r6   r;   r?   r@   rB   rC   rE   rF   rH   rI   rK   rM   rP   rQ   rS   rT   rU   rW   rY   rZ   r[   r^   rŸ   r`   ra   rd   re   rf   rg   rh   ri   rj   rm   rn   ro   rp   rq   rr   rs   rt   ru   rw   rx   ry   rz   r{   r|   r}   r~   r   r€   rƒ   r…   r‰   rŒ   r�   rŽ   r�   r‘   r’   rŒ  Úmainr   r   r   r   Ú<module>   s¨   



     S
