o
    8ήcp                     @   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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.d/ Z%d0d1 Z&d2d3 Z'd4d5 Z(d6d7 Z)d8d9 Z*d:d; Z+d<d= Z,ej-d>d?d@dA Z.ej-d>d?dBdC Z/dDdE Z0dFdG Z1dHdI Z2dJdK Z3dLdM Z4dNdO Z5dPdQ Z6dRdS Z7dTdU Z8dVdW Z9dXdY Z:dZd[ Z;d\d] Z<d^d_ Z=d`da Z>dbdc Z?ddde Z@G dfdg dgeZAeBdhkreC  dS dS )i    N)cudaint64)compile_ptx)f2)unittestCUDATestCaseskip_on_cudasimskip_unless_cc_53c                 C   s   t jj}|| d< d S Nr   r   	threadIdxxaryi r   N/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudapy/test_intrinsics.pysimple_threadidx      r   c                 C   s   t jj}|| |< d S Nr   r   r   r   r   fill_threadidx   r   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 d S )Nr   r   r   r   gridsize)r   r   r   r   r   r   simple_gridsize1d(   s
   

r&   c                 C   sH   t d\}}t d\}}|dkr |dkr"|| d< || d< d S d S d S )Nr"   r   r   r$   )r   r   r   r   r   r   r   r   simple_gridsize2d/   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< q*q"d S r!   )r   r   gridDimr   blockDimr   shaperange)	cstartXstartYgridXgridYheightwidthr   r   r   r   r   intrinsic_forloop_step7   s   
r3   c                 C      t || d< d S r
   )r   popcr   r,   r   r   r   simple_popcB      r7   c                 C   s   t |||| d< d S r
   )r   fmar   abr,   r   r   r   
simple_fmaF      r=   c                 C      t j|d |d | d< d S r
   r   fp16haddr   r;   r<   r   r   r   simple_haddJ      rD   c                 C      t j||| d< d S r
   r@   rC   r   r   r   simple_hadd_scalarN   r>   rG   c                 C   s$   t j|d |d |d | d< d S r
   r   rA   hfmar:   r   r   r   simple_hfmaR   s   $rJ   c                 C   s   t j|||| d< d S r
   rH   r:   r   r   r   simple_hfma_scalarV      rK   c                 C   r?   r
   r   rA   hsubrC   r   r   r   simple_hsubZ   rE   rO   c                 C   rF   r
   rM   rC   r   r   r   simple_hsub_scalar^   r>   rP   c                 C   r?   r
   r   rA   hmulrC   r   r   r   simple_hmulb   rE   rS   c                 C   rF   r
   rQ   rC   r   r   r   simple_hmul_scalarf   r>   rT   c                 C      t j|d | d< d S r
   r   rA   hnegr   r;   r   r   r   simple_hnegj   rL   rY   c                 C      t j|| d< d S r
   rV   rX   r   r   r   simple_hneg_scalarn      r[   c                 C   rU   r
   r   rA   habsrX   r   r   r   simple_habsr   rL   r_   c                 C   rZ   r
   r]   rX   r   r   r   simple_habs_scalarv   r\   r`   c                 C   rF   r
   )r   rA   heqrC   r   r   r   simple_heq_scalarz   r>   rb   c                 C   rF   r
   )r   rA   hnerC   r   r   r   simple_hne_scalar~   r>   rd   c                 C   rF   r
   )r   rA   hgerC   r   r   r   simple_hge_scalar   r>   rf   c                 C   rF   r
   )r   rA   hgtrC   r   r   r   simple_hgt_scalar   r>   rh   c                 C   rF   r
   )r   rA   hlerC   r   r   r   simple_hle_scalar   r>   rj   c                 C   rF   r
   r   rA   hltrC   r   r   r   simple_hlt_scalar   r>   rm   T)devicec                 C      t j| |S r   rk   r   r   r   r   r   
hlt_func_1      rq   c                 C   ro   r   rk   rp   r   r   r   
hlt_func_2   rr   rs   c                 C   s   t ||o	t||| d< d S r
   )rq   rs   rr;   r<   r,   r   r   r   test_multiple_hcmp_1   s   rv   c                 C   s    t ||otj||| d< d S r
   )rq   r   rA   rl   rt   r   r   r   test_multiple_hcmp_2       rw   c                 C   s    t ||otj||| d< d S r
   )rq   r   rA   re   rt   r   r   r   test_multiple_hcmp_3   rx   ry   c                 C   s$   t j||ot j||| d< d S r
   rk   rt   r   r   r   test_multiple_hcmp_4      $rz   c                 C   s$   t j||ot j||| d< d S r
   )r   rA   rl   re   rt   r   r   r   test_multiple_hcmp_5   r{   r|   c                 C   rF   r
   )r   rA   hmaxrC   r   r   r   simple_hmax_scalar   r>   r~   c                 C   rF   r
   )r   rA   hminrC   r   r   r   simple_hmin_scalar   r>   r   c                 C   r4   r
   )r   cbrtrX   r   r   r   simple_cbrt   r8   r   c                 C   r4   r
   )r   brevr6   r   r   r   simple_brev   r8   r   c                 C   r4   r
   )r   clzr6   r   r   r   
simple_clz   r8   r   c                 C   r4   r
   )r   ffsr6   r   r   r   
simple_ffs   r8   r   c                 C   s   t || d< d S r
   roundr6   r   r   r   simple_round   s   r   c                 C   s   t ||| d< d S r
   r   )r   r,   ndigitsr   r   r   simple_round_to   r8   r   c                 C   sJ   t d}| | dkr|d dkr|| | |< d S d| |< d S d| |< d S )Nr      r"   r         r   )r;   r<   r,   r   r   r   r   branching_with_ifs   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)r;   r<   r,   r   innerr   r   r   branching_with_selps   s   
r   c                 C   s   t d}t j| |< d S r   )r   r   laneidr   r   r   r   simple_laneid   s   
r   c                 C   s   t j| d< d S r
   )r   warpsize)r   r   r   r   simple_warpsize   s   r   c                   @   s  e Zd Zdd Zdd Z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d)d* Ze	d$d+d, Zed-d. Zed/d0 Ze	d$d1d2 Zed3d4 Zed5d6 Ze	d$d7d8 Zed9d: Z ed;d< Z!e	d$d=d> Z"ed?d@ Z#edAdB Z$e	d$dCdD Z%edEdF Z&edGdH Z'edIdJ Z(edKdL Z)dMdN Z*dOdP Z+dQdR Z,e	dSdTdU Z-dVdW Z.dXdY Z/dZd[ Z0d\d] Z1e	dSd^d_ Z2d`da Z3dbdc Z4ddde Z5dfdg Z6e	dSdhdi Z7djdk Z8dldm Z9dndo Z:dpdq Z;drds Z<e	dtdudv Z=dwdx Z>dydz Z?e	dtd{d| Z@d}d~ ZAdS )TestCudaIntrinsicc                 C   s@   t dt}tjdtjd}|d | | |d dk d S )Nvoid(int32[:])r   dtyper   r   r   )r   jitr   nponesint32
assertTrueselfcompiledr   r   r   r   test_simple_threadidx      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_threadidx   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   asfortranarrayr   r   r   r   r   r   f_contigous	  s   z<TestCudaIntrinsic.test_fill3d_threadidx.<locals>.f_contigous)r   r   r   )r   r   r   c_resf_resr   r   r   test_fill3d_threadidx   s   
z'TestCudaIntrinsic.test_fill3d_threadidxc                 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nctaid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?q6| 	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*   
fill_valuer   r   r   r   r   r"   z	\s+bra\s+	branching)err_msgr   r   )r   r   r   r   r   r   fullr   r   inspect_asmr   lenrefindalltestingassert_array_equal)
r   sigcu_branching_with_ifscu_branching_with_selpsnr<   r,   expectedr;   ptxr   r   r   	test_selp1  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_gridsize2dJ  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 ]5\}
}||
 || }}t||	|D ]!}t|||D ]}| 	|||f || k|||f || f q\qTqAd S r   )
r   r   r3   r   r   r   r*   zipr+   r   )r   r   r   r   r*   r   r/   r0   r1   r2   r   r   r-   r.   r   r   r   r   r   test_intrinsic_forloop_stepT  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   r;   r<   r,   r   r   r   foof  s   z*TestCudaIntrinsic.test_3dgrid.<locals>.fooi  r   	   )r   r   r   r   )r   r   r   r   r   reshaper   assert_equal)r   r   arrr   r   r   test_3dgride  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o9|t jjt jjt jj  ko9|t jjt jjt jj  k}|t jjt j	j koW|t jjt j	j koW|t jjt j	j k}|o[|| |||f< d S r   )
r   r   r%   r   r   blockIdxr)   r   r   r(   )	r   r   r   r   r;   r<   r,   grid_is_rightgridsize_is_rightr   r   r   r   r  s   z,TestCudaIntrinsic.test_3dgrid_2.<locals>.foo)   r      r   ))r   r   r"   )r   r"   r   )r   r   r   r   bool_r   r   r   )r   r   r   r   r   r   r   r   r   test_3dgrid_2q  s   

"zTestCudaIntrinsic.test_3dgrid_2c                 C   @   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   r7   r   r   r   assertEqualsr   r   r   r   test_popc_u4     zTestCudaIntrinsic.test_popc_u4c                 C   r   )Nzvoid(int32[:], uint64)r   r   r   l        @ r   r   r   r   r   r   r   test_popc_u8  r   zTestCudaIntrinsic.test_popc_u8c                 C   F   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     zTestCudaIntrinsic.test_fma_f4c                 C   r   )
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  r  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 Nzvoid(f2[:], f2[:], f2[:])r   r   r   r   r   r   )	r   r   rD   r   r   float16arrayr   r   r   r   r   arg1arg2r   r   r   	test_hadd     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   rG   r   r   r  r   r   r   r   r   r  r	  refr   r   r   test_hadd_scalar     

z"TestCudaIntrinsic.test_hadd_scalarz(Compilation unsupported in the simulatorc                 C   4   t d d  t t f}tt|dd\}}| d| d S )Nr   r   cczadd.f16)r   r   rG   assertInr   argsr   _r   r   r   test_hadd_ptx     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   rJ   r   r   r  r  r   r   )r   r   r   r  r	  arg3r   r   r   	test_hfma  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   rK   r   r   r  r   r   )r   r   r   r  r	  r  r  r   r   r   test_hfma_scalar  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   rK   r  r  r   r   r   test_hfma_ptx  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   rO   r   r   r  r  r   r   r  r   r   r   	test_hsub  r  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   rP   r   r   r  r   r   r  r   r   r   test_hsub_scalar  r  z"TestCudaIntrinsic.test_hsub_scalarc                 C   r  )Nr  r  zsub.f16)r   r   rP   r  r  r   r   r   test_hsub_ptx  r  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   rS   r   r   r  r  r   r   r  r   r   r   	test_hmul  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   rT   r   r   r  r   r   r  r   r   r   test_hmul_scalar  r  z"TestCudaIntrinsic.test_hmul_scalarc                 C   r  )Nr  r  zmul.f16)r   r   rT   r  r  r   r   r   test_hmul_ptx  r  zTestCudaIntrinsic.test_hmul_ptxc                 C   sV   t dt}tjdtjd}tjdgtjd}|d || tj|d |  d S )Nzvoid(f2[:], f2[:])r   r   r   r   r   )	r   r   rY   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   r[   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   r[   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   r_   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   r`   r   r   r  r.  r   r   r*  r   r   r   test_habs_scalar#  s   
z"TestCudaIntrinsic.test_habs_scalarc                 C   sP   t d d  t f}tt|dd\}}tj dk r | |d d S | d| d S )Nr  r  )r   r"   zand\.b16.*0x7FFF;zabs.f16)r   r   r`   r   runtimeget_versionassertRegexr  r  r   r   r   test_habs_ptx,  s
   zTestCudaIntrinsic.test_habs_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m 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 d    n1 sw   Y  qd S )
N)opzvoid(b1[:], f2, f2)r   r   r"   r   r   r   r   )rb   rd   rf   rh   rj   rm   operatoreqnegegtleltr   subTestr   r   r   r   bool8r  r   )r   fnsopsfnr5  kernelr   gotr	  r  arg4r   r   r   test_fp16_comparison5  s4   





z&TestCudaIntrinsic.test_fp16_comparisonc              	   C   s   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)rA  zvoid(b1[:], f2, f2, f2)r   r   r   r   r   r   r   )rv   rw   ry   rz   r|   r=  r   r   r   r   r>  r  r   )r   	functionsrA  r   r   r  r	  r  r   r   r   !test_multiple_float16_comparisonsU  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   r~   r   r   r  r   r   floatr  r   r   r   	test_hmaxf     


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 rH  )	r   r   r   r   r   r  r   r   rI  r  r   r   r   	test_hminr  rK  zTestCudaIntrinsic.test_hminc                 C   J   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   cbrt_argr   r   r   test_cbrt_f32~  
   zTestCudaIntrinsic.test_cbrt_f32c                 C   rM  )Nzvoid(float64[:], float64)r   r   g      @r   r   rN  )r   r   r   r   r   r  r   r   rO  r   r   r   test_cbrt_f64  rR  zTestCudaIntrinsic.test_cbrt_f64c                 C   r   )Nzvoid(uint32[:], uint32)r   r   r   i0  r   i  )r   r   r   r   r   uint32r   r   r   r   r   test_brev_u4  r   zTestCudaIntrinsic.test_brev_u4z.only get given a Python "int", assumes 32 bitsc                 C   r   )Nzvoid(uint64[:], uint64)r   r   r   l   0  C r   l       `x)r   r   r   r   r   uint64r   r   r   r   r   test_brev_u8     zTestCudaIntrinsic.test_brev_u8c                 C   r   )Nvoid(int32[:], int32)r   r   r      r      r   r   r   r   r   r   r   r   r   r   r   test_clz_i4  r   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   rZ  r   r[  Nr\  r   r   r   r   test_clz_u4  s   	zTestCudaIntrinsic.test_clz_u4c                 C   s@   t dt}tjdtjd}|d |d | |d d d S NrY  r   r   r   l    r   r\  r   r   r   r   test_clz_i4_1s  r   z TestCudaIntrinsic.test_clz_i4_1sc                 C   sB   t dt}tjdtjd}|d |d | |d dd d S )NrY  r   r   r   r   r   CUDA semanticsr\  r   r   r   r   test_clz_i4_0s  s   z TestCudaIntrinsic.test_clz_i4_0sc                 C   r   )Nvoid(int32[:], int64)r   r   r      r   /   r\  r   r   r   r   test_clz_i8  rX  zTestCudaIntrinsic.test_clz_i8c                 C   ^   t dt}tjdtjd}|d |d | |d d |d |d | |d d	 d S )
NrY  r   r   r   rZ  r              r   r   r   r   r   r   r   r   r   r   r   r   test_ffs_i4     zTestCudaIntrinsic.test_ffs_i4c                 C   rg  )
Nr   r   r   r   rZ  r   rh  ri  r   rj  r   r   r   r   test_ffs_u4  rl  zTestCudaIntrinsic.test_ffs_u4c                 C   s@   t dt}tjdtjd}|d |d | |d d d S r_  rj  r   r   r   r   test_ffs_i4_1s  r   z TestCudaIntrinsic.test_ffs_i4_1sc                 C   s@   t dt}tjdtjd}|d |d | |d d d S )NrY  r   r   r   r   rj  r   r   r   r   test_ffs_i4_0s  r   z TestCudaIntrinsic.test_ffs_i4_0sc                 C   rg  )
Nrc  r   r   r   rd  r   r   l        !   rj  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   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   ra  )r   r   r   r   r   r   r   r   r   r   r   test_simple_warpsize  r   z&TestCudaIntrinsic.test_simple_warpsizec                 C   N   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     zTestCudaIntrinsic.test_round_f4c                 C   rv  )Nzvoid(int64[:], float64)r   r   rw  r   r   rx  ry  r   r   r   test_round_f8  r{  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 ]0\}}| j||d |d ||| | j|d	 t||d
d W d    n1 scw   Y  q8d 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   randomseedastypeconcatenater  infnan	itertoolsproductr=  assertPreciseEqualr   r   r   r   valsdigitsr  r   r   r   r   test_round_to_f4  s   "	z"TestCudaIntrinsic.test_round_to_f4z$Overflow behavior differs on CPythonc                 C   T   t dt}tjdtjd}ttjj}d}|d ||| | |d | d S )Nr}  r   r   i,  r   r   )	r   r   r   r   r   r   finfomaxr   r   r   r   r  r   r   r   r   test_round_to_f4_overflow  s   z+TestCudaIntrinsic.test_round_to_f4_overflowc                 C   T   t dt}tjdtjd}d}d}|d ||| | j|d t||dd	 d S )
Nr}  r   r   gQ?r   r   r   r  r  )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   s0  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 ]0\}}| j||d |d ||| | j|d	 t||d
d W d    n1 s_w   Y  q4d}d}| j||d |d ||| | j|d	 t||dd W d    d S 1 sw   Y  d S )N void(float64[:], float64, int32)r   r   r~  r   )r  r  r  r  r  r   r   r"   r   r   r   r  r   r   exactr  g`8p=<   double)r   r   r   r   r   r  r  r  r  r  r  r  r  r  r=  r  r   r  r   r   r   test_round_to_f85  s,   ""z"TestCudaIntrinsic.test_round_to_f8c                 C   r  )Nr  r   r   r   r   r   )	r   r   r   r   r   r  r  r  r   r  r   r   r   test_round_to_f8_overflowL  s   z+TestCudaIntrinsic.test_round_to_f8_overflowc                 C   r  )
Nr  r   r   g\(\?r   r   r   r  r  )r   r   r   r   r   r  r  r   r  r   r   r   test_round_to_f8_halfwayY  s   z*TestCudaIntrinsic.test_round_to_f8_halfwayN)B__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/  r0  r4  rE  rG  rJ  rL  rQ  rS  rU  rW  r]  r^  r`  rb  rf  rk  rm  rn  ro  rq  rt  ru  rz  r|  r  r  r  r  r  r  r   r   r   r   r      s    



	

	




	


	
















r   __main__)Dr  numpyr   r6  r   numbar   r   
numba.cudar   numba.core.typesr   numba.cuda.testingr   r   r   r	   r   r   r   r    r#   r&   r'   r3   r7   r=   rD   rG   rJ   rK   rO   rP   rS   rT   rY   r[   r_   r`   rb   rd   rf   rh   rj   rm   r   rq   rs   rv   rw   ry   rz   r|   r~   r   r   r   r   r   r   r   r   r   r   r   r   r  mainr   r   r   r   <module>   s    



    
x