U
    hâËdê…  ã                   @   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   ú`/home/sam/Atlas/atlas_env/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=   Ú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    ra   c                 C   s   t j | |¡S r   r]   r`   r   r   r   Ú
hlt_func_2¤   s    rb   c                 C   s   t ||ƒot||ƒ| d< d S r   )ra   rb   ©Úrr9   r:   r-   r   r   r   Útest_multiple_hcmp_1©   s    re   c                 C   s    t ||ƒotj ||¡| d< d S r   )ra   r   r=   r^   rc   r   r   r   Útest_multiple_hcmp_2®   s    rf   c                 C   s    t ||ƒotj ||¡| d< d S r   )ra   r   r=   rX   rc   r   r   r   Útest_multiple_hcmp_3³   s    rg   c                 C   s$   t j ||¡ot j ||¡| d< d S r   r]   rc   r   r   r   Útest_multiple_hcmp_4¸   s    rh   c                 C   s$   t j ||¡ot j ||¡| d< d S r   )r   r=   r^   rX   rc   r   r   r   Útest_multiple_hcmp_5½   s    ri   c                 C   s   t j ||¡| d< d S r   )r   r=   Zhmaxr>   r   r   r   Úsimple_hmax_scalarÂ   s    rj   c                 C   s   t j ||¡| d< d S r   )r   r=   Zhminr>   r   r   r   Úsimple_hmin_scalarÆ   s    rk   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    Úlenr=   Zhsin©rd   r   r   r   r   r   Úsimple_hsinÊ   s    
rn   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhcosrm   r   r   r   Úsimple_hcosÑ   s    
ro   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhlogrm   r   r   r   Úsimple_hlogØ   s    
rp   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhlog2rm   r   r   r   Úsimple_hlog2ß   s    
rq   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhlog10rm   r   r   r   Úsimple_hlog10æ   s    
rr   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhexprm   r   r   r   Úsimple_hexpí   s    
rs   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhexp2rm   r   r   r   Úsimple_hexp2ô   s    
rt   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhsqrtrm   r   r   r   Úsimple_hsqrtû   s    
ru   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhrsqrtrm   r   r   r   Úsimple_hrsqrt  s    
rv   c                 C   s   | d S )Ng      à¿r   )r   Údtyper   r   r   Únumpy_hrsqrt
  s    rx   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhceilrm   r   r   r   Úsimple_hceil  s    
ry   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhfloorrm   r   r   r   Úsimple_hfloor  s    
rz   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhrcprm   r   r   r   Úsimple_hrcp  s    
r{   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhtruncrm   r   r   r   Úsimple_htrunc#  s    
r|   c                 C   s.   t  d¡}|t| ƒk r*t j || ¡| |< d S r   )r   r    rl   r=   Zhrintrm   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   ©rw   ©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_valuerw   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È   rl   Ú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§   ))rn   ro   rp   rq   rr   ru   ry   rz   r{   r|   r}   rv   rs   rt   r–   ÚsinÚcosÚlogÚlog2Úlog10ÚsqrtÚceilÚfloorZ
reciprocalÚtruncZrintrx   r«   Zexp2r—   r˜   r  r  rõ   r  rØ   ÚsubTestr   r    rÑ   rï   )rš   ZkernelsZexp_kernelsZexpected_functionsZexpected_exp_functionsrª   r   rd   Ú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    rl   r=   Zhexp10rm   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š   r5  rª   r   rd   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Ø   r2  r   r    r–   r¯   Úbool8rõ   rÈ   )rš   ÚfnsZopsr(  r7  r3  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   )re   rf   rg   rh   ri   r2  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    rj   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 rD  )r   r    rk   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   rG  )r   r    r~   r–   r¯   rñ   rÑ   rï   rH  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Ÿ   rN  r   rO  NrP  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 ©NrM  r   rž   rŸ   l   ÿÿ r   rP  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 )NrM  r   rž   rŸ   r   rË   úCUDA semanticsrP  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   é/   rP  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 )
NrM  r   rž   rŸ   rN  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Ÿ   rN  r   r[  r\  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 rS  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 )NrM  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 )
NrW  r   rž   rŸ   rX  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Ë   rU  )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ž   rg  rŸ   r   rh  ri  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Úproductr2  ÚassertPreciseEqualrƒ   ©rš   r¥   r   ÚvalsÚdigitsrt  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 )Nrl  r   rž   i,  rŸ   r   )	r   r    r†   r–   r¯   rî   ÚfinfoÚmaxrÈ   ©rš   r¥   r   rt  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 )
Nrl  r   rž   gìQ¸…ëÕ?r‰   rŸ   r   ru  rv  )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ž   rm  rË   )rn  ro  rp  rq  rr  r   r   r#   r‰   r‡   r­   rs  rŸ   r   Úexactrv  g`¢8„p=Ö<é   Údouble)r   r    r†   r–   r¯   rñ   r—   r˜   rx  rö   ry  rz  r{  r|  r2  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Š  rv  )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&  r4  r6  rB  rC  rE  rF  rI  rJ  rK  rL  rQ  rR  rT  rV  rZ  r^  r_  r`  ra  rc  re  rf  rj  rk  r�  r…  r†  r‹  rŒ  r�  Ú__classcell__r   r   r›   r   r“   m  sÔ   





	

	




	


	
	






"









r“   Ú__main__)Zr{  Únumpyr–   r8  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    ra   rb   re   rf   rg   rh   ri   rj   rk   rn   ro   rp   rq   rr   rs   rt   ru   rv   rx   ry   rz   r{   r|   r}   r~   r   r€   r�   r„   r†   rŠ   r�   rŽ   r�   r‘   r’   r“   rŽ  Úmainr   r   r   r   Ú<module>   s¨   



     S
