o
    8ήc                     @   sv  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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 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@dfdg ZAdhdi ZBdjdk ZCdldm ZDdndo ZEdpdq ZFdrds ZGdtdu ZHdvdw ZIdxdy ZJdzd{ ZKd|d} ZLd~d ZMdd ZNdd ZOdd ZPdd ZQdd ZRdd ZSdd ZTdd ZUdd ZVdd ZWdd ZXdd ZYdd ZZdd Z[e[d\Z\Z]Z^Z_e[d\Z`ZaZbZce[d\ZdZeZfZge[d\ZhZiZjZkdd ZlG dd deZmendkreo  dS dS )    N)dedent)cudauint32uint64float32float64)unittestCUDATestCaseskip_unless_cc_50cc_X_or_above)configT)devicec                 C      t | S N)r   num r   K/tmp/pip-target-vg8gfxp4/lib/python/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_shared6   s    r4   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$   r1   r'   r2   r3   r*   r   r   r   atomic_binary_2dim_globalF   s   $r6   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_globalP   s
   r7   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_addZ      r?   c              
   C   r8   )Nr.   r:   r   Tr;   r>   r   r   r   atomic_add_wrap_   r@   rA   c                 C      t | dtdtjjtd d S Nr.         Fr4   r   r   r<   r=   r   r>   r   r   r   atomic_add2d      

rH   c                 C   rB   )Nr.   rD   TrG   r>   r   r   r   atomic_add2_wrapi   rI   rJ   c                 C   rB   rC   )r4   r   r   r<   r=   r   r>   r   r   r   atomic_add3n   rI   rK   c              
   C   r8   N      ?r:           Fr+   r   r   r<   r=   r   r>   r   r   r   atomic_add_floats   r@   rP   c              
   C   r8   NrM   r:   rN   TrO   r>   r   r   r   atomic_add_float_wrapx   r@   rR   c                 C   rB   NrM   rD   Fr4   r   r   r<   r=   r   r>   r   r   r   atomic_add_float_2}   rI   rU   c                 C   rB   NrM   rD   TrT   r>   r   r   r   atomic_add_float_2_wrap   rI   rW   c                 C   rB   rS   )r4   r   r   r<   r=   r   r>   r   r   r   atomic_add_float_3   rI   rX   c                 C      t || ddtjjd d S Nr:   rM   Fr7   r   r<   r=   r    r   r   r   r   atomic_add_double_global      r]   c                 C   rY   )Nr:   rM   Tr[   r\   r   r   r   atomic_add_double_global_wrap   r^   r_   c                 C      t | dtjjtd d S Nr.   Fr6   r   r<   r=   r   r>   r   r   r   atomic_add_double_global_2      rc   c                 C   r`   )Nr.   Trb   r>   r   r   r   atomic_add_double_global_2_wrap   rd   re   c                 C   r`   ra   )r6   r   r<   r=   r   r>   r   r   r   atomic_add_double_global_3      rf   c              
   C       t || dtdtjjtdd	 d S rL   r+   r   r   r<   r=   r   r\   r   r   r   atomic_add_double   r@   rj   c              
   C   rh   rQ   ri   r\   r   r   r   atomic_add_double_wrap   r@   rk   c                 C   rB   rS   r4   r   r   r<   r=   r   r>   r   r   r   atomic_add_double_2   rI   rm   c                 C   rB   rV   rl   r>   r   r   r   atomic_add_double_2_wrap   rI   rn   c                 C   rB   rS   )r4   r   r   r<   r=   r   r>   r   r   r   atomic_add_double_3   rI   ro   c              
   C   r8   r9   )r+   r   r   r<   subr   r>   r   r   r   
atomic_sub   r@   rq   c                 C   rB   rC   )r4   r   r   r<   rp   r   r>   r   r   r   atomic_sub2   rI   rr   c                 C   rB   rC   )r4   r   r   r<   rp   r   r>   r   r   r   atomic_sub3   rI   rs   c              
   C   r8   rL   )r+   r   r   r<   rp   r   r>   r   r   r   atomic_sub_float   r@   rt   c                 C   rB   rS   )r4   r   r   r<   rp   r   r>   r   r   r   atomic_sub_float_2   rI   ru   c                 C   rB   rS   )r4   r   r   r<   rp   r   r>   r   r   r   atomic_sub_float_3   rI   rv   c              
   C   rh   rL   )r+   r   r   r<   rp   r   r\   r   r   r   atomic_sub_double   r@   rw   c                 C   rB   rS   )r4   r   r   r<   rp   r   r>   r   r   r   atomic_sub_double_2   rI   rx   c                 C   rB   rS   r4   r   r   r<   rp   r   r>   r   r   r   atomic_sub_double_3   rI   rz   c                 C   rY   rZ   )r7   r   r<   rp   r\   r   r   r   atomic_sub_double_global   r^   r{   c                 C   r`   )NrM   F)r6   r   r<   rp   r   r>   r   r   r   atomic_sub_double_global_2   rg   r|   c                 C   rB   rS   ry   r>   r   r   r   atomic_sub_double_global_3   rI   r}   c              
   C       t | | |tdtjjtdd	 d S )Nr:   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 NrD   F)r4   r   r   r<   r   r   r   r   r   r   atomic_and2   rI   r   c                 C   r   r   )r4   r   r   r<   r   r   r   r   r   r   atomic_and3   rI   r   c                 C      t || d|tjjd d S Nr:   F)r7   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)r6   r   r<   r   r   r   r   r   r   atomic_and_global_2     r   c              
   C   r~   Nr:   r   F)r+   r   r   r<   or_r   r   r   r   r   	atomic_or  r@   r   c                 C   r   r   )r4   r   r   r<   r   r   r   r   r   r   
atomic_or2  rI   r   c                 C   r   r   )r4   r   r   r<   r   r   r   r   r   r   
atomic_or3  rI   r   c                 C   r   r   )r7   r   r<   r   r   r   r   r   atomic_or_global  r^   r   c                 C   r   r   )r6   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   )r4   r   r   r<   r   r   r   r   r   r   atomic_xor2*  rI   r   c                 C   r   r   )r4   r   r   r<   r   r   r   r   r   r   atomic_xor3/  rI   r   c                 C   r   r   )r7   r   r<   r   r   r   r   r   atomic_xor_global4  r^   r   c                 C   r   r   )r6   r   r<   r   r   r   r   r   r   atomic_xor_global_28  r   r   c                 C      t | ||tdtjjt d S Nr:   )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_inc64B  r   r   c                 C   r   r   )r4   r   r   r<   r   r   r   r   r   r   atomic_inc2_32G  rI   r   c                 C   r   r   )r4   r   r   r<   r   r   r   r   r   r   atomic_inc2_64L  rI   r   c                 C   r   r   )r4   r   r   r<   r   r   r   r   r   r   atomic_inc3Q  rI   r   c                 C   r   r   )r7   r   r<   r   r   r   r   r   atomic_inc_globalV  r^   r   c                 C   r   r   )r6   r   r<   r   r   r   r   r   r   atomic_inc_global_2Z  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_dec64d  r   r   c                 C   r   r   )r4   r   r   r<   r   r   r   r   r   r   atomic_dec2_32i  rI   r   c                 C   r   r   )r4   r   r   r<   r   r   r   r   r   r   atomic_dec2_64n  rI   r   c                 C   r   r   )r4   r   r   r<   r   r   r   r   r   r   atomic_dec3s  rI   r   c                 C   r   r   )r7   r   r<   r   r   r   r   r   atomic_dec_globalx  r^   r   c                 C   r   r   )r6   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   )r4   r   r   r<   r   r   r   r   r   r   atomic_exch2  rI   r   c                 C   r   r   )r4   r   r   r<   r   r   r   r   r   r   atomic_exch3  rI   r   c                 C   r   r   )r7   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gidoutr   r   r   atomic_compare_and_swap  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
edd Zdd Zdd Ze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[ Z0d\d] Z1d^d_ Z2d`da Z3dbdc Z4ddde Z5dfdg Z6dhdi Z7djdk Z8dldm Z9dndo Z:dpdq Z;drds Z<dtdu Z=dvdw Z>dxdy Z?dzd{ Z@d|d} ZAd~d 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 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 Zdd Zdd 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   r:   r   zvoid(uint32[:])r.   r:   dtyper.   )r   r   randintastyper   copyr   jitr?   rA   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   r:   r   rE   rF   zvoid(uint32[:,:])r.   rD   r.   )r   r   r   r   r   reshaper   r   r   rH   rJ   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   rK   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   r:   r   zvoid(float32[:])r   r   rM   )r   r   r   r   r   r   intpr   r   rP   rR   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   r:   r   rE   rF   zvoid(float32[:,:])r   r.   )r   r   r   r   r   r   r   r   r   rU   rW   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   rX   r   r   r   r   r   r   test_atomic_add_float_3  r   z'TestCudaAtomics.test_atomic_add_float_3Tc                 C   sp   t jrd S tt|  }tddr&|r| d| d S | d| d S |r0| d| d S | d| d S )N   r   zatom.shared.add.f64zatom.add.f64zatom.shared.cas.b64zatom.cas.b64)r   ENABLE_CUDASIMnextiterinspect_asmvaluesr   assertIn)r   kernelr   asmr   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   r:   r   r   void(int64[:], float64[:])r   r   rM   )r   r   r   int64r   r   r   r   r   rj   rk   r   r   r   testingassert_equalr  )r   r    r   r   cuda_fnwrap_fnr   r   r   r   r   test_atomic_add_double1  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   r:   r   rE   rF   void(float64[:,:])r   r.   )r   r   r   r   r   r   r   r   r   rm   rn   r  r	  r  )r   r   r   r   r
  cuda_fn_wrapr   r   r   test_atomic_add_double_2F  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   ro   r  r	  r  r   r   r   	cuda_funcr   r   r   test_atomic_add_double_3V  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   r:   r  r  r   r   rM   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_global_  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   r:   r   rE   rF   r  r   r.   Fr  )r   r   r   r   r   r   r   r   r   rc   re   r  r	  r  )r   r   r   r   r  r  r  r   r   r   test_atomic_add_double_global_2u  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   rf   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   rq   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   rr   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   rs   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   rt   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   ru   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   rv   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   rw   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   rx   r  r	  r  r   r   r   test_atomic_sub_double_2  
   "z(TestCudaAtomics.test_atomic_sub_double_2c                 C   r*  r  )r   r   r   r   r   r   r   r   r   rz   r  r	  r  r   r   r   test_atomic_sub_double_3  r,  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   r{   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   r*  r  )r   r   r   r   r   r   r   r   r   r|   r  r	  r  r   r   r   test_atomic_sub_double_global_2  r,  z/TestCudaAtomics.test_atomic_sub_double_global_2c                 C   r*  r  )r   r   r   r   r   r   r   r   r   r}   r  r	  r  r   r   r   test_atomic_sub_double_global_3  r,  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   r:   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 	Nr1  r   r:   r   rE   rF   void(uint32[:,:], uint32)r   )r   r   r   r   r   r   r   r   r   r   r   r   r   r4  r   r   cuda_atomic_and2r   r   r   test_atomic_and2     "z TestCudaAtomics.test_atomic_and2c                 C   r6  r7  )r   r   r   r   r   r   r   r   r   r   r   r   r   r4  r   r   cuda_atomic_and3r   r   r   test_atomic_and3  r<  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 Nr1  r   r:   r  zvoid(int32[:], int32[:], int32)r   )r   r   r   int32r   r   r   r   r   r   r  r	  r   r4  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 r7  )r   r   r   r   r   r   r   r   r   r   r  r	  r   r4  r   r   r  r   r   r   test_atomic_and_global_2#     "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 Nr1  r   r:   r   r2  r   r   )r   r   r   r   r   r   r   r   r   r   r   r   r   r   r3  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 r7  )r   r   r   r   r   r   r   r   r   r   r   r   r9  r   r   r   test_atomic_or28  r<  zTestCudaAtomics.test_atomic_or2c                 C   rK  r7  )r   r   r   r   r   r   r   r   r   r   r   r   r=  r   r   r   test_atomic_or3@  r<  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 r@  )r   r   r   rA  r   r   r   r   r   r   r  r	  rB  r   r   r   test_atomic_or_globalH  rD  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 r7  )r   r   r   r   r   r   r   r   r   r   r  r	  rE  r   r   r   test_atomic_or_global_2V  rG  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 rH  )r   r   r   r   r   r   r   r   r   r   r   r   r   r   r3  r   r   r   test_atomic_xor^  rJ  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 r7  )r   r   r   r   r   r   r   r   r   r   r   r   )r   r4  r   r   cuda_atomic_xor2r   r   r   test_atomic_xor2k  r<  z TestCudaAtomics.test_atomic_xor2c                 C   rQ  r7  )r   r   r   r   r   r   r   r   r   r   r   r   )r   r4  r   r   cuda_atomic_xor3r   r   r   test_atomic_xor3s  r<  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 r@  )r   r   r   rA  r   r   r   r   r   r   r  r	  )r   r4  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 r7  )r   r   r   r   r   r   r   r   r   r   r  r	  rE  r   r   r   test_atomic_xor_global_2  rG  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 )Nr:   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 )Nr:   r   r   r   rE   rF   )r   r   r   r   r   )r   r   rY  rZ  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    rY  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-   r^  r`  r   r   r   check_inc_index2  rd  z TestCudaAtomics.check_inc_index2c           	   	   C   sL   |  }t||}|||f || tj|t||kd|d  d S r-   r^  	r   r   rY  r  ra  rb  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.   r:   )r\  r   r   rc  r   r   r4  r   r    r  r   r   r   test_atomic_inc_32     z"TestCudaAtomics.test_atomic_inc_32c              	   C   rh  Nr   z"void(uint64[:], uint64[:], uint64)r.   r:   )r\  r   r   rc  r   rk  r   r   r   test_atomic_inc_64  rm  z"TestCudaAtomics.test_atomic_inc_64c                 C   ,   |  tj\}}d}| |||ddt d S Nr8  r.   rD   )r]  r   r   rg  r   r   r4  r   r  r   r   r   test_atomic_inc2_32     z#TestCudaAtomics.test_atomic_inc2_32c                 C   rp  Nvoid(uint64[:,:], uint64)r.   rD   )r]  r   r   rg  r   rr  r   r   r   test_atomic_inc2_64  rt  z#TestCudaAtomics.test_atomic_inc2_64c                 C   rp  rq  )r]  r   r   rg  r   rr  r   r   r   test_atomic_inc3  rt  z TestCudaAtomics.test_atomic_inc3c              	   C   rh  ri  )r\  r   r   re  r   rk  r   r   r   test_atomic_inc_global_32  
   z)TestCudaAtomics.test_atomic_inc_global_32c              	   C   rh  rn  )r\  r   r   re  r   rk  r   r   r   test_atomic_inc_global_64  rz  z)TestCudaAtomics.test_atomic_inc_global_64c                 C   rp  rq  )r]  r   r   rg  r   rr  r   r   r   test_atomic_inc_global_2_32  rt  z+TestCudaAtomics.test_atomic_inc_global_2_32c                 C   rp  ru  )r]  r   r   rg  r   rr  r   r   r   test_atomic_inc_global_2_64  rt  z+TestCudaAtomics.test_atomic_inc_global_2_64c           
      C   s\   |  }t||}	|	||f ||| tj|t|dk|t||k||d  d S r-   r^  r`  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-   r^  r`  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-   r^  rf  r   r   r   	check_dec  s   

zTestCudaAtomics.check_decc              	   C   rh  ri  )r\  r   r   r~  r   rk  r   r   r   test_atomic_dec_32  rm  z"TestCudaAtomics.test_atomic_dec_32c              	   C   rh  rn  )r\  r   r   r~  r   rk  r   r   r   test_atomic_dec_64  rm  z"TestCudaAtomics.test_atomic_dec_64c                 C   rp  rq  )r]  r   r   r  r   rr  r   r   r   test_atomic_dec2_32  rt  z#TestCudaAtomics.test_atomic_dec2_32c                 C   rp  ru  )r]  r   r   r  r   rr  r   r   r   test_atomic_dec2_64  rt  z#TestCudaAtomics.test_atomic_dec2_64c                 C   rp  rq  )r]  r   r   r  r   rr  r   r   r   test_atomic_dec3_new  rt  z$TestCudaAtomics.test_atomic_dec3_newc              	   C   rh  ri  )r\  r   r   r  r   rk  r   r   r   test_atomic_dec_global_32  rz  z)TestCudaAtomics.test_atomic_dec_global_32c              	   C   rh  rn  )r\  r   r   r  r   rk  r   r   r   test_atomic_dec_global_64  rz  z)TestCudaAtomics.test_atomic_dec_global_64c                 C   rp  rq  )r]  r   r   r  r   rr  r   r   r   test_atomic_dec_global2_32  rt  z*TestCudaAtomics.test_atomic_dec_global2_32c                 C   rp  ru  )r]  r   r   r  r   rr  r   r   r   test_atomic_dec_global2_64"  rt  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   r:   r   rj  r   )r   r   r   r   r   rX  r   r   r   r  r	  )r   r4  r   r    r  r   r   r   test_atomic_exch'  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   r:   r   rE   rF   r8  r   )r   r   r   r   r   r   r   r   r   r  r	  r   r4  r   r  r   r   r   test_atomic_exch21  
   "z!TestCudaAtomics.test_atomic_exch2c                 C   r  )Nr  r  r   r   r:   r   rE   rF   rv  r   )r   r   r   r   r   r   r   r   r   r  r	  r  r   r   r   test_atomic_exch39  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   r:   r   r  rj  r   )
r   r   r   r   rX  r   r   r   r  r	  )r   r4  r    r   r  r  r   r   r   test_atomic_exch_globalA  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r:   r:   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_maxK  s   

z TestCudaAtomics.check_atomic_maxc                 C      | j tjddd d S N   r   r  r  )r  r   rA  r   r   r   r   test_atomic_max_int32S     z%TestCudaAtomics.test_atomic_max_int32c                 C   r  Nr   r  r  )r  r   r   r   r   r   r   test_atomic_max_uint32V  r  z&TestCudaAtomics.test_atomic_max_uint32c                 C   r  r  )r  r   r  r   r   r   r   test_atomic_max_int64Y  r  z%TestCudaAtomics.test_atomic_max_int64c                 C   r  r  )r  r   r   r   r   r   r   test_atomic_max_uint64\  r  z&TestCudaAtomics.test_atomic_max_uint64c                 C   r  r  )r  r   r   r   r   r   r   test_atomic_max_float32_  r  z'TestCudaAtomics.test_atomic_max_float32c                 C   r  r  )r  r   r   r   r   r   r   test_atomic_max_doubleb  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_normalizedindexe     
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      r:   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_oneindexo  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_miny  s   

z TestCudaAtomics.check_atomic_minc                 C   r  r  )r  r   rA  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   r:   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_sharedc                 C   s   |g|d  |g|d   }t j| t j||d}t |}t jjdd|jd|j}||k}||k}	t |}
|| |
|< ||
|	< t |}|| ||< |||	< t	
t}|d |||| t j|
| t j|| d S )N   r   r.   
   r   )r  r  )r   r   shuffleasarray
zeros_liker   r   r   r   r   r   r   r  assert_array_equal)r   nfillunfillr   r   r   r   	fill_maskunfill_mask
expect_res
expect_outr  r   r   r   check_compare_and_swap  s"   



z&TestCudaAtomics.check_compare_and_swapc                 C      | j dddtjd d S )Nr  ir  r  r  r   )r  r   rA  r   r   r   r   test_atomic_compare_and_swap  rd   z,TestCudaAtomics.test_atomic_compare_and_swapc                 C   r  )Nr  ir  r  )r  r   r  r   r   r   r   test_atomic_compare_and_swap2  rd   z-TestCudaAtomics.test_atomic_compare_and_swap2c                 C   @   t jjddt jd}t jjddt jd}| jd||t jd d S Nr  r1  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   test_atomic_compare_and_swap4
  r  z-TestCudaAtomics.test_atomic_compare_and_swap4c                 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_old  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  &  r  zBTestCudaAtomics.test_atomic_max_returns_no_replace.<locals>.kernelr  r  r  r   r   r   "test_atomic_max_returns_no_replace%  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  -  r  zCTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernelr.   r  r  r   r   r   #test_atomic_max_returns_old_replace,  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  4  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_array3     
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  ;     zCTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernelr  r  r  r   r   r   #test_atomic_max_returns_old_nan_val:  r  z3TestCudaAtomics.test_atomic_max_returns_old_nan_valc                 C   r  )Nc                 S   r  Nr      r.   r   r<   r  r  r   r   r   r  B  r  zFTestCudaAtomics.test_atomic_min_returns_old_no_replace.<locals>.kernelr  r  r  r   r   r   &test_atomic_min_returns_old_no_replaceA  r  z6TestCudaAtomics.test_atomic_min_returns_old_no_replacec                 C   r  )Nc                 S   r  r  r
  r  r   r   r   r  I  r  zCTestCudaAtomics.test_atomic_min_returns_old_replace.<locals>.kernelr	  r  r  r   r   r   #test_atomic_min_returns_old_replaceH  r  z3TestCudaAtomics.test_atomic_min_returns_old_replacec                 C   r  )Nc                 S   r  r  r
  r  r   r   r   r  P  r  zHTestCudaAtomics.test_atomic_min_returns_old_nan_in_array.<locals>.kernelr  r  r   r   r   (test_atomic_min_returns_old_nan_in_arrayO  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  W  r  zCTestCudaAtomics.test_atomic_min_returns_old_nan_val.<locals>.kernelr	  r  r  r   r   r   #test_atomic_min_returns_old_nan_valV  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  r  )r  r   rA  r   r   r   r   test_atomic_nanmax_int32i     
z(TestCudaAtomics.test_atomic_nanmax_int32c                 C      | j tjdddd d S Nr   r  r  )r  r   r   r   r   r   r   test_atomic_nanmax_uint32m  r  z)TestCudaAtomics.test_atomic_nanmax_uint32c                 C   r  r  )r  r   r  r   r   r   r   test_atomic_nanmax_int64q  r  z(TestCudaAtomics.test_atomic_nanmax_int64c                 C   r  r  )r  r   r   r   r   r   r   test_atomic_nanmax_uint64u  r  z)TestCudaAtomics.test_atomic_nanmax_uint64c                 C      | j tjddtjd d S Nr  r  r  )r  r   r   r  r   r   r   r   test_atomic_nanmax_float32y     
z*TestCudaAtomics.test_atomic_nanmax_float32c                 C   r  r  )r  r   r   r  r   r   r   r   test_atomic_nanmax_double}  r!  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   r:   r   r.   r  r   r  r   )r   r   r   r   r   r  r   r   r   r   atomic_nanmax_double_sharedr  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  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_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	  r  r   r   r   check_atomic_nanmin  s   

z#TestCudaAtomics.check_atomic_nanminc                 C   r  r  )r+  r   rA  r   r   r   r   test_atomic_nanmin_int32  r  z(TestCudaAtomics.test_atomic_nanmin_int32c                 C   r  r  )r+  r   r   r   r   r   r   test_atomic_nanmin_uint32  r  z)TestCudaAtomics.test_atomic_nanmin_uint32c                 C   r  r  )r+  r   r  r   r   r   r   test_atomic_nanmin_int64  r  z(TestCudaAtomics.test_atomic_nanmin_int64c                 C   r  r  )r+  r   r   r   r   r   r   test_atomic_nanmin_uint64  r  z)TestCudaAtomics.test_atomic_nanmin_uint64c                 C   r  r  )r+  r   r   r  r   r   r   r   test_atomic_nanmin_float  r!  z(TestCudaAtomics.test_atomic_nanmin_floatc                 C   r  r  )r+  r   r   r  r   r   r   r   test_atomic_nanmin_double  r!  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 r#  )r   r   r   r   r   r  r   r   r   r   atomic_nanmin_double_sharedr*  r  r	  r  r   r   r    test_atomic_nanmin_double_shared  r&  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 r'  )r   r   r   r   r   r  r   r   r   r  r*  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<   r  r  r   r   r   r    r  zITestCudaAtomics.test_atomic_nanmax_returns_old_no_replace.<locals>.kernelr  r   r   r6  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  r7  r  r   r   r   r    r  zFTestCudaAtomics.test_atomic_nanmax_returns_old_replace.<locals>.kernelr.   r8  r  r   r   r   &test_atomic_nanmax_returns_old_replace  r  z6TestCudaAtomics.test_atomic_nanmax_returns_old_replacec                 C   r  )Nc                 S   r  r-   r7  r  r   r   r   r    r  zKTestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array.<locals>.kernelr   r   r6  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<   r  r   r  r  r   r   r   r    r  zFTestCudaAtomics.test_atomic_nanmax_returns_old_nan_val.<locals>.kernelr  r8  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<   r*  r  r   r   r   r    r  zITestCudaAtomics.test_atomic_nanmin_returns_old_no_replace.<locals>.kernelr  r8  r  r   r   r   )test_atomic_nanmin_returns_old_no_replace  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	  r8  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    r  zKTestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_array.<locals>.kernelr;  r  r   r   r   +test_atomic_nanmin_returns_old_nan_in_array  r  z;TestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_arrayc                 C   r  )Nc                 S   r  r-   )r   r<   r*  r   r  r  r   r   r   r    r  zFTestCudaAtomics.test_atomic_nanmin_returns_old_nan_val.<locals>.kernelr	  r8  r  r   r   r   &test_atomic_nanmin_returns_old_nan_val  r  z6TestCudaAtomics.test_atomic_nanmin_returns_old_nan_val)T)__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/  r0  r5  r;  r?  rC  rF  rI  rL  rM  rN  rO  rP  rS  rU  rV  rW  r\  r]  rc  re  rg  rl  ro  rs  rw  rx  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-  r.  r/  r0  r1  r3  r4  r6  r9  r:  r<  r=  r?  r@  rA  rB  __classcell__r   r   r   r   r     s   
	
				



	
	


	
	
r   __main__)pnumpyr   textwrapr   numbar   r   r   r   r   numba.cuda.testingr   r	   r
   r   
numba.corer   r   r   r   r   r+   r,   r4   r6   r7   r?   rA   rH   rJ   rK   rP   rR   rU   rW   rX   r]   r_   rc   re   rf   rj   rk   rm   rn   ro   rq   rr   rs   rt   ru   rv   rw   rx   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  r  $atomic_nanmax_double_normalizedindexatomic_nanmax_double_oneindexr$  r)  $atomic_nanmin_double_normalizedindexatomic_nanmin_double_oneindexr2  r   r   rC  mainr   r   r   r   <module>   s    














	
	%        
S