o
    ۾iQ                     @   s  d dl Zd dlmZ d dlmZmZmZmZm	Z	 d dl
mZmZmZ d dlmZ ejdddd	 Zejddd
d Zejdddd Zejdddd Zejdddd Zejdddd Zejdddd Zejd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 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+d>d? Z,d@dA Z-dBdC Z.dDdE Z/dFdG Z0dHdI Z1dJdK Z2dLdM Z3dNdO Z4dPdQ Z5dRdS Z6dTdU Z7dVdW Z8dXdY Z9dZd[ Z:d\d] Z;d^d_ Z<d`da Z=dbdc Z>ddde Z?dfdg Z@dhdi ZAdjdk ZBdldm ZCdndo ZDdpdq ZEdrds ZFdtdu ZGdvdw ZHdxdy ZIdzd{ ZJd|d} ZKd~d ZLdd ZMdd ZNdd ZOdd ZPdd ZQdd ZRdd ZSdd ZTdd ZUdd ZVdd ZWdd ZXdd ZYdd ZZeZd\Z[Z\Z]Z^eZd\Z_Z`ZaZbeZd\ZcZdZeZfeZd\ZgZhZiZjdd Zkdd Zldd ZmG dd deZneodkrep  dS dS )    N)dedent)cudauint32uint64float32float64)unittestCUDATestCasecc_X_or_above)configT)devicec                 C      t | S N)r   num r   X/home/ubuntu/.local/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_atomics.pyatomic_cast_to_uint64	      r   c                 C   r   r   )intr   r   r   r   atomic_cast_to_int   r   r   c                 C   s   | S r   r   r   r   r   r   atomic_cast_none   s   r   c	                 C   sf   t jj}	t j||}
||
|	< t   |||	 | }|r!|| }||
|| t   |
|	 | |	< d S r   r   	threadIdxxsharedarraysyncthreads)aryidxop2	ary_dtypeary_nelements
binop_func	cast_funcinitializerneg_idxtidsmbinr   r   r   atomic_binary_1dim_shared   s   r*   c           
      C   s^   t jj}t j||}| | ||< t   ||| | }	|||	| t   || | |< d S r   r   )
r   r   r    r!   r"   r#   r$   r'   r(   r)   r   r   r   atomic_binary_1dim_shared2(   s   r+   c                 C   s   t jj}t jj}t j||}	| ||f |	||f< t   |||f}
|r5|
d |d  |
d |d  f}
||	|
| t   |	||f | ||f< d S Nr      )r   r   r   yr   r   r   )r   r    r!   	ary_shaper#   y_cast_funcr&   txtyr(   r)   r   r   r   atomic_binary_2dim_shared5   s    r3   c                 C   sT   t jj}t jj}|||f}|r"|d | jd  |d | jd  f}|| || d S r,   )r   r   r   r.   shape)r   r    r#   r0   r&   r1   r2   r)   r   r   r   atomic_binary_2dim_globalE   s   $r5   c                 C   s4   t jj}t|| | }|r|| }|| || d S r   )r   r   r   r   )r   r   r"   r    r#   r&   r'   r)   r   r   r   atomic_binary_1dim_globalO   s
   r6   c              
   C       t | | dtdtjjtdd	 d S Nr-       r   Fr*   r   r   atomicaddr   r   r   r   r   
atomic_addY      r>   c              
   C   r7   )Nr-   r9   r   Tr:   r=   r   r   r   atomic_add_wrap^   r?   r@   c                 C      t | dtdtjjtd d S Nr-         Fr3   r   r   r;   r<   r   r=   r   r   r   atomic_add2c      

rG   c                 C   rA   )Nr-   rC   TrF   r=   r   r   r   atomic_add2_wraph   rH   rI   c                 C   rA   rB   )r3   r   r   r;   r<   r   r=   r   r   r   atomic_add3m   rH   rJ   c              
   C   r7   N      ?r9           Fr*   r   r   r;   r<   r   r=   r   r   r   atomic_add_floatr   r?   rO   c              
   C   r7   NrL   r9   rM   TrN   r=   r   r   r   atomic_add_float_wrapw   r?   rQ   c                 C   rA   NrL   rC   Fr3   r   r   r;   r<   r   r=   r   r   r   atomic_add_float_2|   rH   rT   c                 C   rA   NrL   rC   TrS   r=   r   r   r   atomic_add_float_2_wrap   rH   rV   c                 C   rA   rR   )r3   r   r   r;   r<   r   r=   r   r   r   atomic_add_float_3   rH   rW   c                 C      t || ddtjjd d S Nr9   rL   Fr6   r   r;   r<   r   r   r   r   r   atomic_add_double_global      r\   c                 C   rX   )Nr9   rL   TrZ   r[   r   r   r   atomic_add_double_global_wrap   r]   r^   c                 C      t | dtjjtd d S Nr-   Fr5   r   r;   r<   r   r=   r   r   r   atomic_add_double_global_2      rb   c                 C   r_   )Nr-   Tra   r=   r   r   r   atomic_add_double_global_2_wrap   rc   rd   c                 C   r_   r`   )r5   r   r;   r<   r   r=   r   r   r   atomic_add_double_global_3      re   c              
   C       t || dtdtjjtdd	 d S rK   r*   r   r   r;   r<   r   r[   r   r   r   atomic_add_double   r?   ri   c              
   C   rg   rP   rh   r[   r   r   r   atomic_add_double_wrap   r?   rj   c                 C   rA   rR   r3   r   r   r;   r<   r   r=   r   r   r   atomic_add_double_2   rH   rl   c                 C   rA   rU   rk   r=   r   r   r   atomic_add_double_2_wrap   rH   rm   c                 C   rA   rR   )r3   r   r   r;   r<   r   r=   r   r   r   atomic_add_double_3   rH   rn   c              
   C   r7   r8   )r*   r   r   r;   subr   r=   r   r   r   
atomic_sub   r?   rp   c                 C   rA   rB   )r3   r   r   r;   ro   r   r=   r   r   r   atomic_sub2   rH   rq   c                 C   rA   rB   )r3   r   r   r;   ro   r   r=   r   r   r   atomic_sub3   rH   rr   c              
   C   r7   rK   )r*   r   r   r;   ro   r   r=   r   r   r   atomic_sub_float   r?   rs   c                 C   rA   rR   )r3   r   r   r;   ro   r   r=   r   r   r   atomic_sub_float_2   rH   rt   c                 C   rA   rR   )r3   r   r   r;   ro   r   r=   r   r   r   atomic_sub_float_3   rH   ru   c              
   C   rg   rK   )r*   r   r   r;   ro   r   r[   r   r   r   atomic_sub_double   r?   rv   c                 C   rA   rR   )r3   r   r   r;   ro   r   r=   r   r   r   atomic_sub_double_2   rH   rw   c                 C   rA   rR   r3   r   r   r;   ro   r   r=   r   r   r   atomic_sub_double_3   rH   ry   c                 C   rX   rY   )r6   r   r;   ro   r[   r   r   r   atomic_sub_double_global   r]   rz   c                 C   r_   )NrL   F)r5   r   r;   ro   r   r=   r   r   r   atomic_sub_double_global_2   rf   r{   c                 C   rA   rR   rx   r=   r   r   r   atomic_sub_double_global_3   rH   r|   c              
   C       t | | |tdtjjtdd	 d S )Nr9   r-   F)r*   r   r   r;   and_r   r   r    r   r   r   
atomic_and   r?   r   c                 C      t | |tdtjjtd d S NrC   F)r3   r   r   r;   r~   r   r   r   r   r   atomic_and2   rH   r   c                 C   r   r   )r3   r   r   r;   r~   r   r   r   r   r   atomic_and3   rH   r   c                 C      t || d|tjjd d S Nr9   F)r6   r   r;   r~   r   r   r    r   r   r   atomic_and_global  r]   r   c                 C      t | |tjjtd d S NF)r5   r   r;   r~   r   r   r   r   r   atomic_and_global_2     r   c              
   C   r}   Nr9   r   F)r*   r   r   r;   or_r   r   r   r   r   	atomic_or  r?   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   
atomic_or2  rH   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   
atomic_or3  rH   r   c                 C   r   r   )r6   r   r;   r   r   r   r   r   atomic_or_global  r]   r   c                 C   r   r   )r5   r   r;   r   r   r   r   r   r   atomic_or_global_2  r   r   c              
   C   r}   r   )r*   r   r   r;   xorr   r   r   r   r   
atomic_xor$  r?   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   atomic_xor2)  rH   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   atomic_xor3.  rH   r   c                 C   r   r   )r6   r   r;   r   r   r   r   r   atomic_xor_global3  r]   r   c                 C   r   r   )r5   r   r;   r   r   r   r   r   r   atomic_xor_global_27  r   r   c                 C      t | ||tdtjjt d S Nr9   )r+   r   r   r;   incr   r   r   r    r   r   r   atomic_inc32<     r   c                 C   r   r   )r+   r   r   r;   r   r   r   r   r   r   atomic_inc64A  r   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   atomic_inc2_32F  rH   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   atomic_inc2_64K  rH   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   atomic_inc3P  rH   r   c                 C   r   r   )r6   r   r;   r   r   r   r   r   atomic_inc_globalU  r]   r   c                 C   r   r   )r5   r   r;   r   r   r   r   r   r   atomic_inc_global_2Y  r   r   c                 C   r   r   )r+   r   r   r;   decr   r   r   r   r   atomic_dec32^  r   r   c                 C   r   r   )r+   r   r   r;   r   r   r   r   r   r   atomic_dec64c  r   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   atomic_dec2_32h  rH   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   atomic_dec2_64m  rH   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   atomic_dec3r  rH   r   c                 C   r   r   )r6   r   r;   r   r   r   r   r   atomic_dec_globalw  r]   r   c                 C   r   r   )r5   r   r;   r   r   r   r   r   r   atomic_dec_global_2{  r   r   c                 C   r   r   )r+   r   r   r;   exchr   r   r   r   r   atomic_exch  r   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   atomic_exch2  rH   r   c                 C   r   r   )r3   r   r   r;   r   r   r   r   r   r   atomic_exch3  rH   r   c                 C   r   r   )r6   r   r;   r   r   r   r   r   atomic_exch_global  r]   r   c                 C   sD   t dj| d}i }t|tttd| |d |d |d |d fS )Na  
    def atomic(res, ary):
        tx = cuda.threadIdx.x
        bx = cuda.blockIdx.x
        {func}(res, 0, ary[tx, bx])

    def atomic_double_normalizedindex(res, ary):
        tx = cuda.threadIdx.x
        bx = cuda.blockIdx.x
        {func}(res, 0, ary[tx, uint64(bx)])

    def atomic_double_oneindex(res, ary):
        tx = cuda.threadIdx.x
        {func}(res, 0, ary[tx])

    def atomic_double_shared(res, ary):
        tid = cuda.threadIdx.x
        smary = cuda.shared.array(32, float64)
        smary[tid] = ary[tid]
        smres = cuda.shared.array(1, float64)
        if tid == 0:
            smres[0] = res[0]
        cuda.syncthreads()
        {func}(smres, 0, smary[tid])
        cuda.syncthreads()
        if tid == 0:
            res[0] = smres[0]
    )func)r   r   r   r;   atomic_double_normalizedindexatomic_double_oneindexatomic_double_shared)r   formatexecr   r   r   )r   fnsldr   r   r   gen_atomic_extreme_funcs  s   r   zcuda.atomic.maxzcuda.atomic.minzcuda.atomic.nanmaxzcuda.atomic.nanminc                 C   s<   t d}|| jk rt j| |d  ||| ||< d S d S Nr-   )r   gridsizer;   compare_and_swapresoldr   fill_valgidr   r   r   atomic_compare_and_swap  s   

$r   c                 C   s6   t d}|| jk rt j| |||| ||< d S d S r   )r   r   r   r;   casr   r   r   r   atomic_cas_1dim  s   

r   c                 C   sT   t d}|d | jd k r&|d | jd k r(t j| |||| ||< d S d S d S )N   r   r-   )r   r   r4   r;   r   r   r   r   r   atomic_cas_2dim  s   
$r   c                       s  e Zd Z fddZdd Zdd Zdd Zd	d
 Zdd Zdd Z	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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 d>d? Z!d@dA Z"dBdC Z#dDdE Z$dFdG Z%dHdI Z&dJdK Z'dLdM Z(dNdO Z)dPdQ Z*dRdS Z+dTdU Z,dVdW Z-dXdY Z.dZd[ Z/d\d] Z0d^d_ Z1d`da Z2dbdc Z3ddde Z4dfdg Z5dhdi Z6djdk Z7dldm Z8dndo Z9dpdq Z:drds Z;dtdu Z<dvdw Z=dxdy Z>dzd{ Z?d|d} Z@d~d ZAdd ZBdd ZCdd ZDdd ZEdd ZFdd ZGdd ZHdd ZIdd ZJdd ZKdd ZLdd ZMdd ZNdd ZOdd ZPdd ZQdd ZRdd ZSdd ZTdd ZUdd ZVdd ZWdd ZXdd ZYdd ZZdd Z[dd Z\dd Z]dd Z^dd Z_dd Z`dd Zadd ZbddÄ Zcd*ddƄZdddȄ Zeddʄ Zfdd̄ Zgdd΄ ZhddЄ Zidd҄ ZjddԄ Zkddք Zldd؄ Zmddڄ Zndd܄ Zoddބ Zpdd Zqdd Zrdd Zsdd Ztdd Zudd Zvdd Zwdd Zxdd Zydd Zzdd Z{dd Z|dd Z}dd Z~dd Zdd Zd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  ZS (+  TestCudaAtomicsc                    s   t    tjd d S )Nr   )supersetUpnprandomseedself	__class__r   r   r     s   
zTestCudaAtomics.setUpc                 C   s   t jjddddt j}| }| }tdt}|d | tdt	}|d | t j
dt jd}t|jD ]}|||   d7  < q<| t ||k | t ||k d S Nr   r9   r   zvoid(uint32[:])r-   r9   dtyper-   )r   r   randintastyper   copyr   jitr>   r@   zerosranger   
assertTrueall)r   r   ary_wraporigcuda_atomic_addcuda_atomic_add_wrapgoldir   r   r   test_atomic_add  s   zTestCudaAtomics.test_atomic_addc                 C      t jjddddt jdd}| }| }tdt	}|d | tdt
}|d | | t ||d k | t ||d k d S 	Nr   r9   r   rD   rE   zvoid(uint32[:,:])r-   rC   r-   )r   r   r   r   r   reshaper   r   r   rG   rI   r   r   )r   r   r   r   cuda_atomic_add2cuda_atomic_add2_wrapr   r   r   test_atomic_add2     "z TestCudaAtomics.test_atomic_add2c                 C   `   t jjddddt jdd}| }tdt	}|d | | 
t ||d k d S r   )r   r   r   r   r   r   r   r   r   rJ   r   r   r   r   r   cuda_atomic_add3r   r   r   test_atomic_add3  
   "z TestCudaAtomics.test_atomic_add3c                 C   s   t jjddddt j}| }| t j}tdt	}|d | tdt
}|d | t jdt jd}t|jD ]}|||   d7  < q@| t ||k | t ||k d S Nr   r9   r   zvoid(float32[:])r   r   rL   )r   r   r   r   r   r   intpr   r   rO   rQ   r   r   r   r   r   r   )r   r   r   r   cuda_atomic_add_floatadd_float_wrapr   r   r   r   r   test_atomic_add_float  s   z%TestCudaAtomics.test_atomic_add_floatc                 C   r   	Nr   r9   r   rD   rE   zvoid(float32[:,:])r   r-   )r   r   r   r   r   r   r   r   r   rT   rV   r   r   )r   r   r   r   r   cuda_func_wrapr   r   r   test_atomic_add_float_2  r   z'TestCudaAtomics.test_atomic_add_float_2c                 C   r   r   )r   r   r   r   r   r   r   r   r   rW   r   r   r   r   r   r   test_atomic_add_float_3"  r   z'TestCudaAtomics.test_atomic_add_float_3Tc                 C   s   t jrd S tt|  }tddr2tj	 dkrd}nd}|r'| d}| 
| d| d S |r<| 
d| d S | 
d	| d S )
N   r   )   r-   redatomz.sharedz.add.f64zatom.shared.cas.b64zatom.cas.b64)r   ENABLE_CUDASIMnextiterinspect_asmvaluesr
   r   runtimeget_versionassertIn)r   kernelr   asminstr   r   r   assertCorrectFloat64Atomics*  s   

z+TestCudaAtomics.assertCorrectFloat64Atomicsc                 C   s   t jjdddt jd}t dt j}| }tdt	}|d || tdt
}|d || t jdt jd}t|jD ]}|||   d7  < q?t j|| t j|| | | | | d S Nr   r9   r   r   void(int64[:], float64[:])r   r   rL   )r   r   r   int64r   r   r   r   r   ri   rj   r   r   r   testingassert_equalr  )r   r   r   r   cuda_fnwrap_fnr   r   r   r   r   test_atomic_add_doubleC  s   
z&TestCudaAtomics.test_atomic_add_doublec                 C   s   t jjddddt jdd}| }| }tdt	}|d | tdt
}|d | t j||d  t j||d  | | | | d S 	Nr   r9   r   rD   rE   void(float64[:,:])r   r-   )r   r   r   r   r   r   r   r   r   rl   rm   r  r  r  )r   r   r   r   r  cuda_fn_wrapr   r   r   test_atomic_add_double_2W  s   "
z(TestCudaAtomics.test_atomic_add_double_2c                 C   sd   t jjddddt jdd}| }tdt	}|d | t j
||d  | | d S r  )r   r   r   r   r   r   r   r   r   rn   r  r  r  r   r   r   	cuda_funcr   r   r   test_atomic_add_double_3g  s   "z(TestCudaAtomics.test_atomic_add_double_3c           	      C   s   t jjdddt jd}t dt j}| }d}t|t	}t|t
}|d || |d || t jdt jd}t|jD ]}|||   d7  < qAt j|| t j|| | j|dd	 | j|dd	 d S )
Nr   r9   r  r  r   r   rL   Fr   )r   r   r   r  r   r   r   r   r   r\   r^   r   r   r   r  r  r  )	r   r   r   r   sigr  wrap_cuda_funcr   r   r   r   r   test_atomic_add_double_globalp  s   z-TestCudaAtomics.test_atomic_add_double_globalc                 C   s   t jjddddt jdd}| }| }d}t|t	}t|t
}|d | |d | t j||d  t j||d  | j|d	d
 | j|d	d
 d S Nr   r9   r   rD   rE   r  r   r-   Fr  )r   r   r   r   r   r   r   r   r   rb   rd   r  r  r  )r   r   r   r   r  r  r   r   r   r   test_atomic_add_double_global_2  s   "z/TestCudaAtomics.test_atomic_add_double_global_2c                 C   sh   t jjddddt jdd}| }tdt	}|d | t j
||d  | j|d	d
 d S r"  )r   r   r   r   r   r   r   r   r   re   r  r  r  r  r   r   r   test_atomic_add_double_global_3  s   "z/TestCudaAtomics.test_atomic_add_double_global_3c                 C   s   t jjddddt j}| }tdt}|d | t j	dt jd}t
|jD ]}|||   d8  < q+| t ||k d S r   )r   r   r   r   r   r   r   r   rp   r   r   r   r   r   )r   r   r   cuda_atomic_subr   r   r   r   r   test_atomic_sub  s   zTestCudaAtomics.test_atomic_subc                 C   `   t jjddddt jdd}| }tdt	}|d | | 
t ||d k d S r   )r   r   r   r   r   r   r   r   r   rq   r   r   r   r   r   cuda_atomic_sub2r   r   r   test_atomic_sub2  
   "z TestCudaAtomics.test_atomic_sub2c                 C   r'  r   )r   r   r   r   r   r   r   r   r   rr   r   r   r   r   r   cuda_atomic_sub3r   r   r   test_atomic_sub3  r+  z TestCudaAtomics.test_atomic_sub3c                 C   s   t jjddddt j}| t j}tdt	}|d | t j
dt jd}t|jD ]}|||   d8  < q/| t ||k d S r   )r   r   r   r   r   r   r   r   r   rs   r   r   r   r   r   )r   r   r   cuda_atomic_sub_floatr   r   r   r   r   test_atomic_sub_float  s   z%TestCudaAtomics.test_atomic_sub_floatc                 C   r'  r   )r   r   r   r   r   r   r   r   r   rt   r   r   r(  r   r   r   test_atomic_sub_float_2  r+  z'TestCudaAtomics.test_atomic_sub_float_2c                 C   r'  r   )r   r   r   r   r   r   r   r   r   ru   r   r   r,  r   r   r   test_atomic_sub_float_3  r+  z'TestCudaAtomics.test_atomic_sub_float_3c                 C   s   t jjdddt jd}t dt j}tdt}|d || t jdt jd}t	|j
D ]}|||   d8  < q-t j|| d S r  )r   r   r   r  r   r   r   r   rv   r   r   r  r  )r   r   r   r  r   r   r   r   r   test_atomic_sub_double  s   z&TestCudaAtomics.test_atomic_sub_doublec                 C   Z   t jjddddt jdd}| }tdt	}|d | t j
||d  d S r  )r   r   r   r   r   r   r   r   r   rw   r  r  r  r   r   r   test_atomic_sub_double_2  
   "z(TestCudaAtomics.test_atomic_sub_double_2c                 C   r4  r  )r   r   r   r   r   r   r   r   r   ry   r  r  r  r   r   r   test_atomic_sub_double_3  r6  z(TestCudaAtomics.test_atomic_sub_double_3c                 C   s   t jjdddt jd}t dt j}d}t|t}|d || t jdt jd}t	|j
D ]}|||   d8  < q/t j|| d S r  )r   r   r   r  r   r   r   r   rz   r   r   r  r  )r   r   r   r  r  r   r   r   r   r   test_atomic_sub_double_global  s   z-TestCudaAtomics.test_atomic_sub_double_globalc                 C   r4  r  )r   r   r   r   r   r   r   r   r   r{   r  r  r  r   r   r   test_atomic_sub_double_global_2  r6  z/TestCudaAtomics.test_atomic_sub_double_global_2c                 C   r4  r  )r   r   r   r   r   r   r   r   r   r|   r  r  r  r   r   r   test_atomic_sub_double_global_3  r6  z/TestCudaAtomics.test_atomic_sub_double_global_3c                 C   s   t jd}t jjddddt j}| }tdt}|d || | }t	|j
D ]}|||   |M  < q.| t ||k d S )N  r   r9   r   void(uint32[:], uint32)r   )r   r   r   r   r   r   r   r   r   r   r   r   r   r   
rand_constr   r   r  r   r   r   r   r   test_atomic_and  s   zTestCudaAtomics.test_atomic_andc                 C   n   t jd}t jjddddt jdd}| }tdt	}|d || | 
t |||@ k d S 	Nr;  r   r9   r   rD   rE   void(uint32[:,:], uint32)r   )r   r   r   r   r   r   r   r   r   r   r   r   r   r>  r   r   cuda_atomic_and2r   r   r   test_atomic_and2     "z TestCudaAtomics.test_atomic_and2c                 C   r@  rA  )r   r   r   r   r   r   r   r   r   r   r   r   r   r>  r   r   cuda_atomic_and3r   r   r   test_atomic_and3  rF  z TestCudaAtomics.test_atomic_and3c                 C   s   t jd}t jjdddt jd}t jjdddt jd}d}t|t}|d ||| | }t|j	D ]}|||   |M  < q6t j
|| d S Nr;  r   r9   r  zvoid(int32[:], int32[:], int32)r   )r   r   r   int32r   r   r   r   r   r   r  r  r   r>  r   r   r  r  r   r   r   r   r   test_atomic_and_global%     z&TestCudaAtomics.test_atomic_and_globalc                 C   sh   t jd}t jjddddt jdd}| }tdt	}|d || t j
|||@  d S rA  )r   r   r   r   r   r   r   r   r   r   r  r  r   r>  r   r   r  r   r   r   test_atomic_and_global_23     "z(TestCudaAtomics.test_atomic_and_global_2c                 C   s   t jd}t jjddddt j}| }tdt}|d || t j	dt jd}t
|jD ]}|||   |O  < q2| t ||k d S Nr;  r   r9   r   r<  r   r   )r   r   r   r   r   r   r   r   r   r   r   r   r   r   r=  r   r   r   test_atomic_or;     zTestCudaAtomics.test_atomic_orc                 C   n   t jd}t jjddddt jdd}| }tdt	}|d || | 
t |||B k d S rA  )r   r   r   r   r   r   r   r   r   r   r   r   rC  r   r   r   test_atomic_or2H  rF  zTestCudaAtomics.test_atomic_or2c                 C   rU  rA  )r   r   r   r   r   r   r   r   r   r   r   r   rG  r   r   r   test_atomic_or3P  rF  zTestCudaAtomics.test_atomic_or3c                 C   s   t jd}t jjdddt jd}t jjdddt jd}d}t|t}|d ||| | }t|j	D ]}|||   |O  < q6t j
|| d S rJ  )r   r   r   rK  r   r   r   r   r   r   r  r  rL  r   r   r   test_atomic_or_globalX  rN  z%TestCudaAtomics.test_atomic_or_globalc                 C   sh   t jd}t jjddddt jdd}| }tdt	}|d || t j
|||B  d S rA  )r   r   r   r   r   r   r   r   r   r   r  r  rO  r   r   r   test_atomic_or_global_2f  rQ  z'TestCudaAtomics.test_atomic_or_global_2c                 C   s   t jd}t jjddddt j}| }tdt}|d || t j	dt jd}t
|jD ]}|||   |N  < q2| t ||k d S rR  )r   r   r   r   r   r   r   r   r   r   r   r   r   r   r=  r   r   r   test_atomic_xorn  rT  zTestCudaAtomics.test_atomic_xorc                 C   n   t jd}t jjddddt jdd}| }tdt	}|d || | 
t |||A k d S rA  )r   r   r   r   r   r   r   r   r   r   r   r   )r   r>  r   r   cuda_atomic_xor2r   r   r   test_atomic_xor2{  rF  z TestCudaAtomics.test_atomic_xor2c                 C   r[  rA  )r   r   r   r   r   r   r   r   r   r   r   r   )r   r>  r   r   cuda_atomic_xor3r   r   r   test_atomic_xor3  rF  z TestCudaAtomics.test_atomic_xor3c                 C   s   t jd}t jjdddt jd}t jjdddt jd}| }d}t|t}|d ||| t|j	D ]}|||   |N  < q6t j
|| d S rJ  )r   r   r   rK  r   r   r   r   r   r   r  r  )r   r>  r   r   r   r  r  r   r   r   r   test_atomic_xor_global  s   z&TestCudaAtomics.test_atomic_xor_globalc                 C   sh   t jd}t jjddddt jdd}| }tdt	}|d || t j
|||A  d S rA  )r   r   r   r   r   r   r   r   r   r   r  r  rO  r   r   r   test_atomic_xor_global_2  rQ  z(TestCudaAtomics.test_atomic_xor_global_2c                 C   s@   t jjd|d}t jjdddd|}t jd|d}|||fS )Nr9   r   r   r   )r   r   r   r   arange)r   r   rconstraryary_idxr   r   r   inc_dec_1dim_setup  s   
z"TestCudaAtomics.inc_dec_1dim_setupc                 C   s8   t jjd|d}t jjdddd|dd}||fS )Nr9   r   r   r   rD   rE   )r   r   r   r   r   )r   r   rc  rd  r   r   r   inc_dec_2dim_setup  s    z"TestCudaAtomics.inc_dec_2dim_setupc           
   	   C   sN   |  }t||}	|	||f ||| tj|t||kd|d  d S r,   r   r   r   r   r  r  where
r   r   r   rc  r  nblocksblksizer   r   r  r   r   r   check_inc_index     $zTestCudaAtomics.check_inc_indexc           
   	   C   sN   |  }t||}	|	||f ||| tj|t||kd|d  d S r,   rh  rj  r   r   r   check_inc_index2  rn  z TestCudaAtomics.check_inc_index2c           	   	   C   sL   |  }t||}|||f || tj|t||kd|d  d S r,   rh  	r   r   rc  r  rk  rl  r   r   r  r   r   r   	check_inc  s   $zTestCudaAtomics.check_incc              	   C   2   | j tjd\}}}d}| ||||ddt d S Nr   "void(uint32[:], uint32[:], uint32)r-   r9   )rf  r   r   rm  r   r   r>  r   r   r  r   r   r   test_atomic_inc_32     z"TestCudaAtomics.test_atomic_inc_32c              	   C   rr  Nr   z"void(uint64[:], uint64[:], uint64)r-   r9   )rf  r   r   rm  r   ru  r   r   r   test_atomic_inc_64  rw  z"TestCudaAtomics.test_atomic_inc_64c                 C   ,   |  tj\}}d}| |||ddt d S NrB  r-   rC   )rg  r   r   rq  r   r   r>  r   r  r   r   r   test_atomic_inc2_32     z#TestCudaAtomics.test_atomic_inc2_32c                 C   rz  Nvoid(uint64[:,:], uint64)r-   rC   )rg  r   r   rq  r   r|  r   r   r   test_atomic_inc2_64  r~  z#TestCudaAtomics.test_atomic_inc2_64c                 C   rz  r{  )rg  r   r   rq  r   r|  r   r   r   test_atomic_inc3  r~  z TestCudaAtomics.test_atomic_inc3c              	   C   rr  rs  )rf  r   r   ro  r   ru  r   r   r   test_atomic_inc_global_32  
   z)TestCudaAtomics.test_atomic_inc_global_32c              	   C   rr  rx  )rf  r   r   ro  r   ru  r   r   r   test_atomic_inc_global_64  r  z)TestCudaAtomics.test_atomic_inc_global_64c                 C   rz  r{  )rg  r   r   rq  r   r|  r   r   r   test_atomic_inc_global_2_32  r~  z+TestCudaAtomics.test_atomic_inc_global_2_32c                 C   rz  r  )rg  r   r   rq  r   r|  r   r   r   test_atomic_inc_global_2_64  r~  z+TestCudaAtomics.test_atomic_inc_global_2_64c           
      C   s\   |  }t||}	|	||f ||| tj|t|dk|t||k||d  d S r,   rh  rj  r   r   r   check_dec_index     

zTestCudaAtomics.check_dec_indexc           
      C   s\   |  }t||}	|	||f ||| tj|t|dk|t||k||d  d S r,   rh  rj  r   r   r   check_dec_index2  r  z TestCudaAtomics.check_dec_index2c           	      C   sZ   |  }t||}|||f || tj|t|dk|t||k||d  d S r,   rh  rp  r   r   r   	check_dec  s   

zTestCudaAtomics.check_decc              	   C   rr  rs  )rf  r   r   r  r   ru  r   r   r   test_atomic_dec_32  rw  z"TestCudaAtomics.test_atomic_dec_32c              	   C   rr  rx  )rf  r   r   r  r   ru  r   r   r   test_atomic_dec_64  rw  z"TestCudaAtomics.test_atomic_dec_64c                 C   rz  r{  )rg  r   r   r  r   r|  r   r   r   test_atomic_dec2_32  r~  z#TestCudaAtomics.test_atomic_dec2_32c                 C   rz  r  )rg  r   r   r  r   r|  r   r   r   test_atomic_dec2_64  r~  z#TestCudaAtomics.test_atomic_dec2_64c                 C   rz  r{  )rg  r   r   r  r   r|  r   r   r   test_atomic_dec3_new  r~  z$TestCudaAtomics.test_atomic_dec3_newc              	   C   rr  rs  )rf  r   r   r  r   ru  r   r   r   test_atomic_dec_global_32!  r  z)TestCudaAtomics.test_atomic_dec_global_32c              	   C   rr  rx  )rf  r   r   r  r   ru  r   r   r   test_atomic_dec_global_64'  r  z)TestCudaAtomics.test_atomic_dec_global_64c                 C   rz  r{  )rg  r   r   r  r   r|  r   r   r   test_atomic_dec_global2_32-  r~  z*TestCudaAtomics.test_atomic_dec_global2_32c                 C   rz  r  )rg  r   r   r  r   r|  r   r   r   test_atomic_dec_global2_642  r~  z*TestCudaAtomics.test_atomic_dec_global2_64c                 C   sn   t jjddt jd}t jjddddt j}t jdt jd}tdt}|d ||| t j	
|| d S )	N2   d   r   r   r9   r   rt  r   )r   r   r   r   r   rb  r   r   r   r  r  )r   r>  r   r   r  r   r   r   test_atomic_exch7  s   z TestCudaAtomics.test_atomic_exchc                 C   d   t jjddt jd}t jjddddt jdd}td	t}|d
 || t j	
|| d S )Nr  r  r   r   r9   r   rD   rE   rB  r   )r   r   r   r   r   r   r   r   r   r  r  r   r>  r   r  r   r   r   test_atomic_exch2A  
   "z!TestCudaAtomics.test_atomic_exch2c                 C   r  )Nr  r  r   r   r9   r   rD   rE   r  r   )r   r   r   r   r   r   r   r   r   r  r  r  r   r   r   test_atomic_exch3I  r  z!TestCudaAtomics.test_atomic_exch3c                 C   sn   t jjddt jd}t jdt jd}t jjdddt jd}d}t|t}|d ||| t j	|| d S )	Nr  r  r   r9   r   r  rt  r   )
r   r   r   r   rb  r   r   r   r  r  )r   r>  r   r   r  r  r   r   r   test_atomic_exch_globalQ  s   z'TestCudaAtomics.test_atomic_exch_globalc                 C   s\   t jj||dd|}t jd|jd}tt}|d || t 	|}t j
|| d S )Nr9   r9   r   r-   r   )r   r   r   r   r   r   r   r   
atomic_maxmaxr  r  r   r   lohivalsr   r  r   r   r   r   check_atomic_max[  s   

z TestCudaAtomics.check_atomic_maxc                 C      | j tjddd d S N   r   r  r  )r  r   rK  r   r   r   r   test_atomic_max_int32c     z%TestCudaAtomics.test_atomic_max_int32c                 C   r  Nr   r  r  )r  r   r   r   r   r   r   test_atomic_max_uint32f  r  z&TestCudaAtomics.test_atomic_max_uint32c                 C   r  r  )r  r   r  r   r   r   r   test_atomic_max_int64i  r  z%TestCudaAtomics.test_atomic_max_int64c                 C   r  r  )r  r   r   r   r   r   r   test_atomic_max_uint64l  r  z&TestCudaAtomics.test_atomic_max_uint64c                 C   r  r  )r  r   r   r   r   r   r   test_atomic_max_float32o  r  z'TestCudaAtomics.test_atomic_max_float32c                 C   r  r  )r  r   r   r   r   r   r   test_atomic_max_doubler  r  z&TestCudaAtomics.test_atomic_max_doublec                 C   s`   t jjddddt j}t dt j}tdt}|d || t 	|}t j
|| d S Nr   r  r  r   r-   void(float64[:], float64[:,:]))r   r   r   r   r   r   r   r   !atomic_max_double_normalizedindexr  r  r  r   r  r   r  r   r   r   r   &test_atomic_max_double_normalizedindexu     
z6TestCudaAtomics.test_atomic_max_double_normalizedindexc                 C   s`   t jjddddt j}t dt j}tdt}|d || t 	|}t j
|| d S Nr      r9   r   r-   void(float64[:], float64[:])r   )r   r   r   r   r   r   r   r   atomic_max_double_oneindexr  r  r  r  r   r   r   test_atomic_max_double_oneindex  r  z/TestCudaAtomics.test_atomic_max_double_oneindexc                 C   s^   t jj||dd|}t jdg|jd}tt}|d || t 	|}t j
|| d S )Nr  r   r  r   )r   r   r   r   r   r   r   r   
atomic_minminr  r  r  r   r   r   check_atomic_min  s   

z TestCudaAtomics.check_atomic_minc                 C   r  r  )r  r   rK  r   r   r   r   test_atomic_min_int32  r  z%TestCudaAtomics.test_atomic_min_int32c                 C   r  r  )r  r   r   r   r   r   r   test_atomic_min_uint32  r  z&TestCudaAtomics.test_atomic_min_uint32c                 C   r  r  )r  r   r  r   r   r   r   test_atomic_min_int64  r  z%TestCudaAtomics.test_atomic_min_int64c                 C   r  r  )r  r   r   r   r   r   r   test_atomic_min_uint64  r  z&TestCudaAtomics.test_atomic_min_uint64c                 C   r  r  )r  r   r   r   r   r   r   test_atomic_min_float  r  z%TestCudaAtomics.test_atomic_min_floatc                 C   r  r  )r  r   r   r   r   r   r   test_atomic_min_double  r  z&TestCudaAtomics.test_atomic_min_doublec                 C   sd   t jjddddt j}t dt jd }tdt}|d || t 	|}t j
|| d S r  )r   r   r   r   r   onesr   r   !atomic_min_double_normalizedindexr  r  r  r  r   r   r   &test_atomic_min_double_normalizedindex     
z6TestCudaAtomics.test_atomic_min_double_normalizedindexc                 C   sd   t jjddddt j}t dt jd }tdt}|d || t 	|}t j
|| d S r  )r   r   r   r   r   r  r   r   atomic_min_double_oneindexr  r  r  r  r   r   r   test_atomic_min_double_oneindex  r  z/TestCudaAtomics.test_atomic_min_double_oneindexc                 C   s`   t d|}tjjddddtj}tdtjtj }|d || tj	
|tjg d S )Nr  r   r  r-   r-   r   r-   )r   r   r   r   r   r   r   r   nanr  r  )r   r   r  r  r   r   r   r    _test_atomic_minmax_nan_location  s
   z0TestCudaAtomics._test_atomic_minmax_nan_locationc                 C   sd   t d|}tjjddddtj}| }tdtjtj	 }|d || tj
|| d S )Nr  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   _test_atomic_minmax_nan_val  s   z+TestCudaAtomics._test_atomic_minmax_nan_valc                 C      |  t d S r   )r  r  r   r   r   r   test_atomic_min_nan_location     z,TestCudaAtomics.test_atomic_min_nan_locationc                 C   r  r   )r  r  r   r   r   r   test_atomic_max_nan_location  r  z,TestCudaAtomics.test_atomic_max_nan_locationc                 C   r  r   )r  r  r   r   r   r   test_atomic_min_nan_val  r  z'TestCudaAtomics.test_atomic_min_nan_valc                 C   r  r   )r  r  r   r   r   r   test_atomic_max_nan_val  r  z'TestCudaAtomics.test_atomic_max_nan_valc                 C   sd   t jjddddt j}t dt j}d}t|t}|d || t 	|}t j
|| d S Nr   r9   r   r-   r  r   )r   r   r   r   r   r   r   r   atomic_max_double_sharedr  r  r  r   r  r   r  r  r   r   r   r   test_atomic_max_double_shared  s   
z-TestCudaAtomics.test_atomic_max_double_sharedc                 C   sh   t jjddddt j}t dt jd }d}t|t}|d || t 	|}t j
|| d S r  )r   r   r   r   r   r  r   r   atomic_min_double_sharedr  r  r  r  r   r   r   test_atomic_min_double_shared  s   
z-TestCudaAtomics.test_atomic_min_double_sharedr-   c                 C   s   |g|d  |g|d   }t j| t j||d}|dkr"d|_t |}t jjdd|jd|j}	||k}
||k}t |}|	|
 ||
< |||< |	 }t
|}|dkrc|d |||	| n	|d |||	| t j|| t j|| d S )	Nr   r   )
   r-   r  r   r  r  )r  r  )r   r   shuffleasarrayr4   
zeros_liker   r   r   r   r   r   r  assert_array_equal)r   nfillunfillr   cas_funcndimr   outr   	fill_maskunfill_mask
expect_res
expect_outr  r   r   r   	check_cas  s&   


zTestCudaAtomics.check_casc                 C      | j dddtjtd d S Nr  r  r  r  r  r   r  )r  r   rK  r   r   r   r   r   test_atomic_compare_and_swap     
z,TestCudaAtomics.test_atomic_compare_and_swapc                 C   r  Nr  r  r  )r  r   r  r   r   r   r   r   test_atomic_compare_and_swap2  r  z-TestCudaAtomics.test_atomic_compare_and_swap2c                 C   B   t jjddt jd}t jjddt jd}| jd||t jtd d S Nr  r;  r   r-      r  r  )r   r   r   r   r  r   r   rfillrunfillr   r   r   test_atomic_compare_and_swap3  
   
z-TestCudaAtomics.test_atomic_compare_and_swap3c                 C   r  r  )r   r   r   r   r  r   r  r   r   r   test_atomic_compare_and_swap4  r  z-TestCudaAtomics.test_atomic_compare_and_swap4c                 C   r  r  )r  r   rK  r   r   r   r   r   test_atomic_cas_1dim%  r  z$TestCudaAtomics.test_atomic_cas_1dimc                 C      | j dddtjtdd d S )Nr  r  r  r   r  r  r  r   r  r  )r  r   rK  r   r   r   r   r   test_atomic_cas_2dim)     
z$TestCudaAtomics.test_atomic_cas_2dimc                 C   r  r  )r  r   r  r   r   r   r   r   test_atomic_cas2_1dim-  r  z%TestCudaAtomics.test_atomic_cas2_1dimc                 C   r  )Nr  r  r  r   r  )r  r   r  r   r   r   r   r   test_atomic_cas2_2dim1  r  z%TestCudaAtomics.test_atomic_cas2_2dimc                 C   r  r  )r   r   r   r   r  r   r  r   r   r   test_atomic_cas3_1dim5  r  z%TestCudaAtomics.test_atomic_cas3_1dimc                 C   D   t jjddt jd}t jjddt jd}| jd||t jtdd d S 	Nr  r;  r   r-   r  r  r   r  )r   r   r   r   r  r   r  r   r   r   test_atomic_cas3_2dim;  
   
z%TestCudaAtomics.test_atomic_cas3_2dimc                 C   r  r  )r   r   r   r   r  r   r  r   r   r   test_atomic_cas4_1dimA  r  z%TestCudaAtomics.test_atomic_cas4_1dimc                 C   r	  r
  )r   r   r   r   r  r   r  r   r   r   test_atomic_cas4_2dimG  r  z%TestCudaAtomics.test_atomic_cas4_2dimc                 C   sZ   t jdt jd}||d< |d | t |r#| t |d  d S | |d | d S )Nr   r   r   r  r-   )r   r   r   isnanr   assertEqualr   r
  initialr   r   r   r   _test_atomic_returns_oldR  s   
z(TestCudaAtomics._test_atomic_returns_oldc                 C      t jdd }| |d d S )Nc                 S      t j| dd| d< d S r,   )r   r;   r<   r   r   r   r   r
  \     z;TestCudaAtomics.test_atomic_add_returns_old.<locals>.kernelr  r   r   r  r   r
  r   r   r   test_atomic_add_returns_old[     
z+TestCudaAtomics.test_atomic_add_returns_oldc                 C   r  )Nc                 S   r  r,   r   r;   r  r  r   r   r   r
  c  r  zBTestCudaAtomics.test_atomic_max_returns_no_replace.<locals>.kernelr  r  r  r   r   r   "test_atomic_max_returns_no_replaceb  r  z2TestCudaAtomics.test_atomic_max_returns_no_replacec                 C   r  )Nc                 S      t j| dd| d< d S Nr   r  r-   r  r  r   r   r   r
  j  r  zCTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernelr-   r  r  r   r   r   #test_atomic_max_returns_old_replacei  r  z3TestCudaAtomics.test_atomic_max_returns_old_replacec                 C       t jdd }| |tj d S )Nc                 S   r  r,   r  r  r   r   r   r
  q  r  zHTestCudaAtomics.test_atomic_max_returns_old_nan_in_array.<locals>.kernelr   r   r  r   r  r  r   r   r   (test_atomic_max_returns_old_nan_in_arrayp     
z8TestCudaAtomics.test_atomic_max_returns_old_nan_in_arrayc                 C   r  )Nc                 S      t j| dtj| d< d S r,   )r   r;   r  r   r  r  r   r   r   r
  x     zCTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernelr  r  r  r   r   r   #test_atomic_max_returns_old_nan_valw  r  z3TestCudaAtomics.test_atomic_max_returns_old_nan_valc                 C   r  )Nc                 S   r  Nr      r-   r   r;   r  r  r   r   r   r
    r  zFTestCudaAtomics.test_atomic_min_returns_old_no_replace.<locals>.kernelr  r  r  r   r   r   &test_atomic_min_returns_old_no_replace~  r  z6TestCudaAtomics.test_atomic_min_returns_old_no_replacec                 C   r  )Nc                 S   r  r  r*  r  r   r   r   r
    r  zCTestCudaAtomics.test_atomic_min_returns_old_replace.<locals>.kernelr)  r  r  r   r   r   #test_atomic_min_returns_old_replace  r  z3TestCudaAtomics.test_atomic_min_returns_old_replacec                 C   r!  )Nc                 S   r  r(  r*  r  r   r   r   r
    r  zHTestCudaAtomics.test_atomic_min_returns_old_nan_in_array.<locals>.kernelr"  r  r   r   r   (test_atomic_min_returns_old_nan_in_array  r$  z8TestCudaAtomics.test_atomic_min_returns_old_nan_in_arrayc                 C   r  )Nc                 S   r%  r,   )r   r;   r  r   r  r  r   r   r   r
    r&  zCTestCudaAtomics.test_atomic_min_returns_old_nan_val.<locals>.kernelr)  r  r  r   r   r   #test_atomic_min_returns_old_nan_val  r  z3TestCudaAtomics.test_atomic_min_returns_old_nan_valc           	      C   sj   t jj||dd|}||dd d< t jd|jd}tt}|d || t 	|}t j
|| d S )Nr  r   r-   r   r   )r   r   r   r   r   r   r   r   atomic_nanmaxnanmaxr  r  	r   r   r  r  init_valr  r   r  r   r   r   r   check_atomic_nanmax  s   

z#TestCudaAtomics.check_atomic_nanmaxc                 C      | j tjdddd d S Nr  r  r   r   r  r  r2  )r3  r   rK  r   r   r   r   test_atomic_nanmax_int32     
z(TestCudaAtomics.test_atomic_nanmax_int32c                 C      | j tjdddd d S Nr   r  r6  )r3  r   r   r   r   r   r   test_atomic_nanmax_uint32  r8  z)TestCudaAtomics.test_atomic_nanmax_uint32c                 C   r4  r5  )r3  r   r  r   r   r   r   test_atomic_nanmax_int64  r8  z(TestCudaAtomics.test_atomic_nanmax_int64c                 C   r9  r:  )r3  r   r   r   r   r   r   test_atomic_nanmax_uint64  r8  z)TestCudaAtomics.test_atomic_nanmax_uint64c                 C      | j tjddtjd d S Nr  r  r6  )r3  r   r   r  r   r   r   r   test_atomic_nanmax_float32     
z*TestCudaAtomics.test_atomic_nanmax_float32c                 C   r>  r?  )r3  r   r   r  r   r   r   r   test_atomic_nanmax_double  rA  z)TestCudaAtomics.test_atomic_nanmax_doublec                 C   sx   t jjddddt j}t j|dd d< t jdg|jd}d}t	|t
}|d || t |}t j|| d S 	Nr   r9   r   r-   r   r   r  r   )r   r   r   r   r   r  r   r   r   r   atomic_nanmax_double_sharedr0  r  r  r  r   r   r    test_atomic_nanmax_double_shared     
z0TestCudaAtomics.test_atomic_nanmax_double_sharedc                 C   sp   t jjddddt j}t j|dd d< t dt j}tdt	}|d || t 
|}t j|| d S 	Nr   r  r9   r   r-   r   r  r   )r   r   r   r   r   r  r   r   r   r  r0  r  r  r  r   r   r   "test_atomic_nanmax_double_oneindex  s   
z2TestCudaAtomics.test_atomic_nanmax_double_oneindexc           	      C   sl   t jj||dd|}||dd d< t jdg|jd}tt}|d || t 	|}t j
|| d S )Nr  r   r-   r   r  r   )r   r   r   r   r   r   r   r   atomic_nanminnanminr  r  r1  r   r   r   check_atomic_nanmin  s   

z#TestCudaAtomics.check_atomic_nanminc                 C   r4  r5  )rK  r   rK  r   r   r   r   test_atomic_nanmin_int32  r8  z(TestCudaAtomics.test_atomic_nanmin_int32c                 C   r9  r:  )rK  r   r   r   r   r   r   test_atomic_nanmin_uint32  r8  z)TestCudaAtomics.test_atomic_nanmin_uint32c                 C   r4  r5  )rK  r   r  r   r   r   r   test_atomic_nanmin_int64  r8  z(TestCudaAtomics.test_atomic_nanmin_int64c                 C   r9  r:  )rK  r   r   r   r   r   r   test_atomic_nanmin_uint64  r8  z)TestCudaAtomics.test_atomic_nanmin_uint64c                 C   r>  r?  )rK  r   r   r  r   r   r   r   test_atomic_nanmin_float  rA  z(TestCudaAtomics.test_atomic_nanmin_floatc                 C   r>  r?  )rK  r   r   r  r   r   r   r   test_atomic_nanmin_double  rA  z)TestCudaAtomics.test_atomic_nanmin_doublec                 C   sx   t jjddddt j}t j|dd d< t jdg|jd}d}t	|t
}|d || t |}t j|| d S rC  )r   r   r   r   r   r  r   r   r   r   atomic_nanmin_double_sharedrJ  r  r  r  r   r   r    test_atomic_nanmin_double_shared  rF  z0TestCudaAtomics.test_atomic_nanmin_double_sharedc                 C   sr   t jjddddt j}t j|dd d< t dgt j}tdt	}|d || t 
|}t j|| d S rG  )r   r   r   r   r   r  r   r   r   r  rJ  r  r  r  r   r   r   "test_atomic_nanmin_double_oneindex  s   
z2TestCudaAtomics.test_atomic_nanmin_double_oneindexc                 C   sx   t jdt jd}||d< t j|d< |d | t |r2| t |d  | t |d  d S | |d | d S )Nr   r   r   r-   r  )r   r   r   r  r  assertFalser   r  r  r   r   r   _test_atomic_nan_returns_old  s   

z,TestCudaAtomics._test_atomic_nan_returns_oldc                 C   r  )Nc                 S   r  r,   r   r;   r0  r  r   r   r   r
    r  zITestCudaAtomics.test_atomic_nanmax_returns_old_no_replace.<locals>.kernelr  r   r   rV  r  r   r   r   )test_atomic_nanmax_returns_old_no_replace  r  z9TestCudaAtomics.test_atomic_nanmax_returns_old_no_replacec                 C   r  )Nc                 S   r  r  rW  r  r   r   r   r
  "  r  zFTestCudaAtomics.test_atomic_nanmax_returns_old_replace.<locals>.kernelr-   rX  r  r   r   r   &test_atomic_nanmax_returns_old_replace!  r  z6TestCudaAtomics.test_atomic_nanmax_returns_old_replacec                 C   r!  )Nc                 S   r  r,   rW  r  r   r   r   r
  )  r  zKTestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array.<locals>.kernelr   r   rV  r   r  r  r   r   r   +test_atomic_nanmax_returns_old_nan_in_array(  r$  z;TestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_arrayc                 C   r  )Nc                 S   r%  r,   )r   r;   r0  r   r  r  r   r   r   r
  0  r&  zFTestCudaAtomics.test_atomic_nanmax_returns_old_nan_val.<locals>.kernelr  rX  r  r   r   r   &test_atomic_nanmax_returns_old_nan_val/  r  z6TestCudaAtomics.test_atomic_nanmax_returns_old_nan_valc                 C   r  )Nc                 S   r  r(  r   r;   rJ  r  r   r   r   r
  7  r  zITestCudaAtomics.test_atomic_nanmin_returns_old_no_replace.<locals>.kernelr  rX  r  r   r   r   )test_atomic_nanmin_returns_old_no_replace6  r  z9TestCudaAtomics.test_atomic_nanmin_returns_old_no_replacec                 C   r  )Nc                 S   r  r  r^  r  r   r   r   r
  >  r  zFTestCudaAtomics.test_atomic_nanmin_returns_old_replace.<locals>.kernelr)  rX  r  r   r   r   &test_atomic_nanmin_returns_old_replace=  r  z6TestCudaAtomics.test_atomic_nanmin_returns_old_replacec                 C   r!  )Nc                 S   r  r(  r^  r  r   r   r   r
  E  r  zKTestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_array.<locals>.kernelr[  r  r   r   r   +test_atomic_nanmin_returns_old_nan_in_arrayD  r$  z;TestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_arrayc                 C   r  )Nc                 S   r%  r,   )r   r;   rJ  r   r  r  r   r   r   r
  L  r&  zFTestCudaAtomics.test_atomic_nanmin_returns_old_nan_val.<locals>.kernelr)  rX  r  r   r   r   &test_atomic_nanmin_returns_old_nan_valK  r  z6TestCudaAtomics.test_atomic_nanmin_returns_old_nan_val)T)r-   )__name__
__module____qualname__r   r   r   r   r   r   r   r  r  r  r  r!  r#  r$  r&  r*  r.  r0  r1  r2  r3  r5  r7  r8  r9  r:  r?  rE  rI  rM  rP  rS  rV  rW  rX  rY  rZ  r]  r_  r`  ra  rf  rg  rm  ro  rq  rv  ry  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  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  r7  r;  r<  r=  r@  rB  rE  rH  rK  rL  rM  rN  rO  rP  rQ  rS  rT  rV  rY  rZ  r\  r]  r_  r`  ra  rb  __classcell__r   r   r   r   r     s(   					



	
	


	
	

r   __main__)qnumpyr   textwrapr   numbar   r   r   r   r   numba.cuda.testingr   r	   r
   
numba.corer   r   r   r   r   r*   r+   r3   r5   r6   r>   r@   rG   rI   rJ   rO   rQ   rT   rV   rW   r\   r^   rb   rd   re   ri   rj   rl   rm   rn   rp   rq   rr   rs   rt   ru   rv   rw   ry   rz   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  r  r  r  r/  $atomic_nanmax_double_normalizedindexatomic_nanmax_double_oneindexrD  rI  $atomic_nanmin_double_normalizedindexatomic_nanmin_double_oneindexrR  r   r   r   r   rc  mainr   r   r   r   <module>   s    














	
	%         
