U
    ,d                     @   sr  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rneo  dS )    N)dedent)cudauint32uint64float32float64)unittestCUDATestCaseskip_unless_cc_50cc_X_or_above)configT)Zdevicec                 C   s   t | S N)r   num r   H/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/tests/cudapy/test_atomics.pyatomic_cast_to_uint64
   s    r   c                 C   s   t | S r   )intr   r   r   r   atomic_cast_to_int   s    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   |||	 | }|rB|| }||
|| t   |
|	 | |	< d S r   r   	threadIdxxsharedarraysyncthreads)aryidxop2	ary_dtypeary_nelements
binop_func	cast_funcZ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}
|rj|
d |d  |
d |d  f}
||	|
| t   |	||f | ||f< d S Nr      )r   r   r   yr   r   r   )r   r   r   Z	ary_shaper!   y_cast_funcr#   txtyr%   r&   r   r   r   atomic_binary_2dim_shared6   s     r/   c                 C   sT   t jj}t jj}|||f}|rD|d | jd  |d | jd  f}|| || d S r)   )r   r   r   r+   shape)r   r   r!   r,   r#   r-   r.   r&   r   r   r   atomic_binary_2dim_globalF   s    $r1   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
    r2   c              
   C   s    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   s       r9   c              
   C   s    t | | dtdtjjtdd	 d S )Nr*   r4   r   Tr5   r8   r   r   r   atomic_add_wrap_   s       r:   c                 C   s   t | dtdtjjtd d S Nr*         Fr/   r   r   r6   r7   r   r8   r   r   r   atomic_add2d   s
    
  r@   c                 C   s   t | dtdtjjtd d S )Nr*   r<   Tr?   r8   r   r   r   atomic_add2_wrapi   s
    
  rA   c                 C   s   t | dtdtjjtd d S r;   )r/   r   r   r6   r7   r   r8   r   r   r   atomic_add3n   s
    
  rB   c              
   C   s    t | | dtdtjjtdd	 d S N      ?r4           Fr'   r   r   r6   r7   r   r8   r   r   r   atomic_add_floats   s       rG   c              
   C   s    t | | dtdtjjtdd	 d S NrD   r4   rE   TrF   r8   r   r   r   atomic_add_float_wrapx   s       rI   c                 C   s   t | dtdtjjtd d S NrD   r<   Fr/   r   r   r6   r7   r   r8   r   r   r   atomic_add_float_2}   s
    
  rL   c                 C   s   t | dtdtjjtd d S NrD   r<   TrK   r8   r   r   r   atomic_add_float_2_wrap   s
    
  rN   c                 C   s   t | dtdtjjtd d S rJ   )r/   r   r   r6   r7   r   r8   r   r   r   atomic_add_float_3   s
    
  rO   c                 C   s   t || ddtjjd d S Nr4   rD   Fr2   r   r6   r7   r   r   r   r   r   atomic_add_double_global   s    rS   c                 C   s   t || ddtjjd d S )Nr4   rD   TrQ   rR   r   r   r   atomic_add_double_global_wrap   s    rT   c                 C   s   t | dtjjtd d S Nr*   Fr1   r   r6   r7   r   r8   r   r   r   atomic_add_double_global_2   s    rW   c                 C   s   t | dtjjtd d S )Nr*   TrV   r8   r   r   r   atomic_add_double_global_2_wrap   s    rX   c                 C   s   t | dtjjtd d S rU   )r1   r   r6   r7   r   r8   r   r   r   atomic_add_double_global_3   s    rY   c              
   C   s    t || dtdtjjtdd	 d S rC   r'   r   r   r6   r7   r   rR   r   r   r   atomic_add_double   s       r[   c              
   C   s    t || dtdtjjtdd	 d S rH   rZ   rR   r   r   r   atomic_add_double_wrap   s       r\   c                 C   s   t | dtdtjjtd d S rJ   r/   r   r   r6   r7   r   r8   r   r   r   atomic_add_double_2   s
    
  r^   c                 C   s   t | dtdtjjtd d S rM   r]   r8   r   r   r   atomic_add_double_2_wrap   s
    
  r_   c                 C   s   t | dtdtjjtd d S rJ   )r/   r   r   r6   r7   r   r8   r   r   r   atomic_add_double_3   s
    
  r`   c              
   C   s    t | | dtdtjjtdd	 d S r3   )r'   r   r   r6   subr   r8   r   r   r   
atomic_sub   s       rb   c                 C   s   t | dtdtjjtd d S r;   )r/   r   r   r6   ra   r   r8   r   r   r   atomic_sub2   s
    
  rc   c                 C   s   t | dtdtjjtd d S r;   )r/   r   r   r6   ra   r   r8   r   r   r   atomic_sub3   s
    
  rd   c              
   C   s    t | | dtdtjjtdd	 d S rC   )r'   r   r   r6   ra   r   r8   r   r   r   atomic_sub_float   s       re   c                 C   s   t | dtdtjjtd d S rJ   )r/   r   r   r6   ra   r   r8   r   r   r   atomic_sub_float_2   s
    
  rf   c                 C   s   t | dtdtjjtd d S rJ   )r/   r   r   r6   ra   r   r8   r   r   r   atomic_sub_float_3   s
    
  rg   c              
   C   s    t || dtdtjjtdd	 d S rC   )r'   r   r   r6   ra   r   rR   r   r   r   atomic_sub_double   s       rh   c                 C   s   t | dtdtjjtd d S rJ   )r/   r   r   r6   ra   r   r8   r   r   r   atomic_sub_double_2   s
    
  ri   c                 C   s   t | dtdtjjtd d S rJ   r/   r   r   r6   ra   r   r8   r   r   r   atomic_sub_double_3   s
    
  rk   c                 C   s   t || ddtjjd d S rP   )r2   r   r6   ra   rR   r   r   r   atomic_sub_double_global   s    rl   c                 C   s   t | dtjjtd d S )NrD   F)r1   r   r6   ra   r   r8   r   r   r   atomic_sub_double_global_2   s    rm   c                 C   s   t | dtdtjjtd d S rJ   rj   r8   r   r   r   atomic_sub_double_global_3   s
    
  rn   c              
   C   s    t | | |tdtjjtdd	 d S )Nr4   r*   F)r'   r   r   r6   and_r   r   r   r   r   r   
atomic_and   s       rq   c                 C   s   t | |tdtjjtd d S Nr<   F)r/   r   r   r6   ro   r   rp   r   r   r   atomic_and2   s
    
  rs   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   ro   r   rp   r   r   r   atomic_and3   s
    
  rt   c                 C   s   t || d|tjjd d S Nr4   F)r2   r   r6   ro   r   r   r   r   r   r   atomic_and_global  s    rw   c                 C   s   t | |tjjtd d S NF)r1   r   r6   ro   r   rp   r   r   r   atomic_and_global_2  s     ry   c              
   C   s    t | | |tdtjjtdd	 d S Nr4   r   F)r'   r   r   r6   or_r   rp   r   r   r   	atomic_or  s       r|   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r{   r   rp   r   r   r   
atomic_or2  s
    
  r}   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r{   r   rp   r   r   r   
atomic_or3  s
    
  r~   c                 C   s   t || d|tjjd d S ru   )r2   r   r6   r{   rv   r   r   r   atomic_or_global  s    r   c                 C   s   t | |tjjtd d S rx   )r1   r   r6   r{   r   rp   r   r   r   atomic_or_global_2   s     r   c              
   C   s    t | | |tdtjjtdd	 d S rz   )r'   r   r   r6   xorr   rp   r   r   r   
atomic_xor%  s       r   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r   r   rp   r   r   r   atomic_xor2*  s
    
  r   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r   r   rp   r   r   r   atomic_xor3/  s
    
  r   c                 C   s   t || d|tjjd d S ru   )r2   r   r6   r   rv   r   r   r   atomic_xor_global4  s    r   c                 C   s   t | |tjjtd d S rx   )r1   r   r6   r   r   rp   r   r   r   atomic_xor_global_28  s     r   c                 C   s   t | ||tdtjjt d S Nr4   )r(   r   r   r6   incr   r   r   r   r   r   r   atomic_inc32=  s     r   c                 C   s   t | ||tdtjjt d S r   )r(   r   r   r6   r   r   r   r   r   r   atomic_inc64B  s     r   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r   r   rp   r   r   r   atomic_inc2_32G  s
    
  r   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r   r   rp   r   r   r   atomic_inc2_64L  s
    
  r   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r   r   rp   r   r   r   atomic_inc3Q  s
    
  r   c                 C   s   t || d|tjjd d S ru   )r2   r   r6   r   rv   r   r   r   atomic_inc_globalV  s    r   c                 C   s   t | |tjjtd d S rx   )r1   r   r6   r   r   rp   r   r   r   atomic_inc_global_2Z  s     r   c                 C   s   t | ||tdtjjt d S r   )r(   r   r   r6   decr   r   r   r   r   atomic_dec32_  s     r   c                 C   s   t | ||tdtjjt d S r   )r(   r   r   r6   r   r   r   r   r   r   atomic_dec64d  s     r   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r   r   rp   r   r   r   atomic_dec2_32i  s
    
  r   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r   r   rp   r   r   r   atomic_dec2_64n  s
    
  r   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r   r   rp   r   r   r   atomic_dec3s  s
    
  r   c                 C   s   t || d|tjjd d S ru   )r2   r   r6   r   rv   r   r   r   atomic_dec_globalx  s    r   c                 C   s   t | |tjjtd d S rx   )r1   r   r6   r   r   rp   r   r   r   atomic_dec_global_2|  s     r   c                 C   s   t | ||tdtjjt d S r   )r(   r   r   r6   exchr   r   r   r   r   atomic_exch  s     r   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r   r   rp   r   r   r   atomic_exch2  s
    
  r   c                 C   s   t | |tdtjjtd d S rr   )r/   r   r   r6   r   r   rp   r   r   r   atomic_exch3  s
    
  r   c                 C   s   t || d|tjjd d S ru   )r2   r   r6   r   rv   r   r   r   atomic_exch_global  s    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   r6   Zatomic_double_normalizedindexZatomic_double_oneindexZatomic_double_shared)r   formatexecr   r   r   )r   fnsZ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 r8t j| |d  ||| }|||< d S )Nr*   )r   Zgridsizer6   Zcompare_and_swap)resoldr   Z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  < qx| t ||k | t ||k d S Nr   r4   r   zvoid(uint32[:])r*   r4   dtyper*   )r   r   randintastyper   copyr   jitr9   r:   zerosranger   
assertTrueall)r   r   ary_wraporigZcuda_atomic_addZcuda_atomic_add_wrapgoldir   r   r   test_atomic_add  s    zTestCudaAtomics.test_atomic_addc                 C   s   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   r4   r   r=   r>   zvoid(uint32[:,:])r*   r<   r*   )r   r   r   r   r   reshaper   r   r   r@   rA   r   r   )r   r   r   r   cuda_atomic_add2Zcuda_atomic_add2_wrapr   r   r   test_atomic_add2  s    "z TestCudaAtomics.test_atomic_add2c                 C   s`   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   rB   r   r   r   r   r   Zcuda_atomic_add3r   r   r   test_atomic_add3  s
    "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   r4   r   zvoid(float32[:])r   r   rD   )r   r   r   r   r   r   intpr   r   rG   rI   r   r   r   r   r   r   )r   r   r   r   Zcuda_atomic_add_floatZadd_float_wrapr   r   r   r   r   test_atomic_add_float  s    z%TestCudaAtomics.test_atomic_add_floatc                 C   s   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   r4   r   r=   r>   zvoid(float32[:,:])r   r*   )r   r   r   r   r   r   r   r   r   rL   rN   r   r   )r   r   r   r   r   Zcuda_func_wrapr   r   r   test_atomic_add_float_2
  s    "z'TestCudaAtomics.test_atomic_add_float_2c                 C   s`   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   rO   r   r   r   r   r   r   test_atomic_add_float_3  s
    "z'TestCudaAtomics.test_atomic_add_float_3Tc                 C   sj   t jr
d S tt|  }tddrH|r:| d| qf| d| n|rZ| d| n| d| d S )N   r   zatom.shared.add.f64zatom.add.f64zatom.shared.cas.b64zatom.cas.b64)r   ZENABLE_CUDASIMnextiterZinspect_asmvaluesr   ZassertIn)r   kernelr   Z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   r4   r   r   void(int64[:], float64[:])r   r   rD   )r   r   r   int64r   r   r   r   r   r[   r\   r   r   r   testingassert_equalr   )r   r   r   r   cuda_fnZ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   r4   r   r=   r>   void(float64[:,:])r   r*   )r   r   r   r   r   r   r   r   r   r^   r_   r   r   r   )r   r   r   r   r   Z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   r`   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  < qt j|| t j|| | j|dd	 | j|dd	 d S )
Nr   r4   r   r   r   r   rD   Fr   )r   r   r   r   r   r   r   r   r   rS   rT   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   r4   r   r=   r>   r   r   r*   Fr   )r   r   r   r   r   r   r   r   r   rW   rX   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   rY   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  < qV| t ||k d S r   )r   r   r   r   r   r   r   r   rb   r   r   r   r   r   )r   r   r   Zcuda_atomic_subr   r   r   r   r   test_atomic_sub  s    zTestCudaAtomics.test_atomic_subc                 C   s`   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   rc   r   r   r   r   r   Zcuda_atomic_sub2r   r   r   test_atomic_sub2  s
    "z TestCudaAtomics.test_atomic_sub2c                 C   s`   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   rd   r   r   r   r   r   Zcuda_atomic_sub3r   r   r   test_atomic_sub3  s
    "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   re   r   r   r   r   r   )r   r   r   Zcuda_atomic_sub_floatr   r   r   r   r   test_atomic_sub_float  s    z%TestCudaAtomics.test_atomic_sub_floatc                 C   s`   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   rf   r   r   r   r   r   r   test_atomic_sub_float_2  s
    "z'TestCudaAtomics.test_atomic_sub_float_2c                 C   s`   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   rg   r   r   r   r   r   r   test_atomic_sub_float_3  s
    "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  < qZt j|| d S r   )r   r   r   r   r   r   r   r   rh   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   sZ   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   ri   r   r   r   r   r   r   test_atomic_sub_double_2  s
    "z(TestCudaAtomics.test_atomic_sub_double_2c                 C   sZ   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   rk   r   r   r   r   r   r   test_atomic_sub_double_3  s
    "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   rl   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   sZ   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   rm   r   r   r   r   r   r   test_atomic_sub_double_global_2  s
    "z/TestCudaAtomics.test_atomic_sub_double_global_2c                 C   sZ   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   test_atomic_sub_double_global_3  s
    "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   r4   r   void(uint32[:], uint32)r   )r   r   r   r   r   r   r   r   rq   r   r   r   r   r   
rand_constr   r   r   r   r   r   r   r   test_atomic_and  s    zTestCudaAtomics.test_atomic_andc                 C   sn   t jd}t jjddddt jdd}| }tdt	}|d || | 
t |||@ k d S 	Nr   r   r4   r   r=   r>   void(uint32[:,:], uint32)r   )r   r   r   r   r   r   r   r   r   rs   r   r   r   r   r   r   Zcuda_atomic_and2r   r   r   test_atomic_and2  s    "z TestCudaAtomics.test_atomic_and2c                 C   sn   t jd}t jjddddt jdd}| }tdt	}|d || | 
t |||@ k d S r   )r   r   r   r   r   r   r   r   r   rt   r   r   r   r   r   r   Zcuda_atomic_and3r   r   r   test_atomic_and3  s    "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  < qlt j
|| d S Nr   r   r4   r   zvoid(int32[:], int32[:], int32)r   )r   r   r   int32r   r   rw   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   test_atomic_and_global  s    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 r   )r   r   r   r   r   r   r   r   r   ry   r   r   r   r   r   r   r   r   r   r   test_atomic_and_global_2#  s    "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  < qd| t ||k d S Nr   r   r4   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+  s    zTestCudaAtomics.test_atomic_orc                 C   sn   t jd}t jjddddt jdd}| }tdt	}|d || | 
t |||B k d S r   )r   r   r   r   r   r   r   r   r   r}   r   r   r  r   r   r   test_atomic_or28  s    "zTestCudaAtomics.test_atomic_or2c                 C   sn   t jd}t jjddddt jdd}| }tdt	}|d || | 
t |||B k d S r   )r   r   r   r   r   r   r   r   r   r~   r   r   r  r   r   r   test_atomic_or3@  s    "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  < qlt j
|| d S r  )r   r   r   r  r   r   r   r   r   r   r   r   r  r   r   r   test_atomic_or_globalH  s    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 r   )r   r   r   r   r   r   r   r   r   r   r   r   r
  r   r   r   test_atomic_or_global_2V  s    "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  < qd| t ||k d S r  )r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   test_atomic_xor^  s    zTestCudaAtomics.test_atomic_xorc                 C   sn   t jd}t jjddddt jdd}| }tdt	}|d || | 
t |||A k d S r   )r   r   r   r   r   r   r   r   r   r   r   r   )r   r   r   r   Zcuda_atomic_xor2r   r   r   test_atomic_xor2k  s    "z TestCudaAtomics.test_atomic_xor2c                 C   sn   t jd}t jjddddt jdd}| }tdt	}|d || | 
t |||A k d S r   )r   r   r   r   r   r   r   r   r   r   r   r   )r   r   r   r   Zcuda_atomic_xor3r   r   r   test_atomic_xor3s  s    "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  < qlt 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_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 r   )r   r   r   r   r   r   r   r   r   r   r   r   r
  r   r   r   test_atomic_xor_global_2  s    "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 )Nr4   r   r   r   )r   r   r   r   arange)r   r   rconstraryZ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 )Nr4   r   r   r   r=   r>   )r   r   r   r   r   )r   r   r  r  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   r  r   nblocksblksizer   r   r   r   r   r   check_inc_index  s    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  s    z TestCudaAtomics.check_inc_index2c           	   	   C   sL   |  }t||}|||f || tj|t||kd|d  d S r)   r  	r   r   r  r   r  r   r   r   r   r   r   r   	check_inc  s    zTestCudaAtomics.check_incc              	   C   s2   | j tjd\}}}d}| ||||ddt d S Nr   "void(uint32[:], uint32[:], uint32)r*   r4   )r  r   r   r!  r   r   r   r   r   r   r   r   r   test_atomic_inc_32  s    z"TestCudaAtomics.test_atomic_inc_32c              	   C   s2   | j tjd\}}}d}| ||||ddt d S Nr   z"void(uint64[:], uint64[:], uint64)r*   r4   )r  r   r   r!  r   r'  r   r   r   test_atomic_inc_64  s    z"TestCudaAtomics.test_atomic_inc_64c                 C   s,   |  tj\}}d}| |||ddt d S Nr  r*   r<   )r  r   r   r$  r   r   r   r   r   r   r   r   test_atomic_inc2_32  s    z#TestCudaAtomics.test_atomic_inc2_32c                 C   s,   |  tj\}}d}| |||ddt d S Nvoid(uint64[:,:], uint64)r*   r<   )r  r   r   r$  r   r,  r   r   r   test_atomic_inc2_64  s    z#TestCudaAtomics.test_atomic_inc2_64c                 C   s,   |  tj\}}d}| |||ddt d S r+  )r  r   r   r$  r   r,  r   r   r   test_atomic_inc3  s    z TestCudaAtomics.test_atomic_inc3c              	   C   s2   | j tjd\}}}d}| ||||ddt d S r%  )r  r   r   r"  r   r'  r   r   r   test_atomic_inc_global_32  s
    z)TestCudaAtomics.test_atomic_inc_global_32c              	   C   s2   | j tjd\}}}d}| ||||ddt d S r)  )r  r   r   r"  r   r'  r   r   r   test_atomic_inc_global_64  s
    z)TestCudaAtomics.test_atomic_inc_global_64c                 C   s,   |  tj\}}d}| |||ddt d S r+  )r  r   r   r$  r   r,  r   r   r   test_atomic_inc_global_2_32  s    z+TestCudaAtomics.test_atomic_inc_global_2_32c                 C   s,   |  tj\}}d}| |||ddt d S r.  )r  r   r   r$  r   r,  r   r   r   test_atomic_inc_global_2_64  s    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  s    
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  s    
z TestCudaAtomics.check_dec_index2c           	      C   sZ   |  }t||}|||f || tj|t|dk|t||k||d  d S r)   r  r#  r   r   r   	check_dec  s    
zTestCudaAtomics.check_decc              	   C   s2   | j tjd\}}}d}| ||||ddt d S r%  )r  r   r   r6  r   r'  r   r   r   test_atomic_dec_32  s    z"TestCudaAtomics.test_atomic_dec_32c              	   C   s2   | j tjd\}}}d}| ||||ddt d S r)  )r  r   r   r6  r   r'  r   r   r   test_atomic_dec_64  s    z"TestCudaAtomics.test_atomic_dec_64c                 C   s,   |  tj\}}d}| |||ddt d S r+  )r  r   r   r8  r   r,  r   r   r   test_atomic_dec2_32  s    z#TestCudaAtomics.test_atomic_dec2_32c                 C   s,   |  tj\}}d}| |||ddt d S r.  )r  r   r   r8  r   r,  r   r   r   test_atomic_dec2_64  s    z#TestCudaAtomics.test_atomic_dec2_64c                 C   s,   |  tj\}}d}| |||ddt d S r+  )r  r   r   r8  r   r,  r   r   r   test_atomic_dec3_new  s    z$TestCudaAtomics.test_atomic_dec3_newc              	   C   s2   | j tjd\}}}d}| ||||ddt d S r%  )r  r   r   r7  r   r'  r   r   r   test_atomic_dec_global_32  s
    z)TestCudaAtomics.test_atomic_dec_global_32c              	   C   s2   | j tjd\}}}d}| ||||ddt d S r)  )r  r   r   r7  r   r'  r   r   r   test_atomic_dec_global_64  s
    z)TestCudaAtomics.test_atomic_dec_global_64c                 C   s,   |  tj\}}d}| |||ddt d S r+  )r  r   r   r8  r   r,  r   r   r   test_atomic_dec_global2_32  s    z*TestCudaAtomics.test_atomic_dec_global2_32c                 C   s,   |  tj\}}d}| |||ddt d S r.  )r  r   r   r8  r   r,  r   r   r   test_atomic_dec_global2_64"  s    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   r4   r   r&  r   )r   r   r   r   r   r  r   r   r   r   r   )r   r   r   r   r   r   r   r   test_atomic_exch'  s    z TestCudaAtomics.test_atomic_exchc                 C   sd   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 )NrB  rC  r   r   r4   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_exch21  s
    "z!TestCudaAtomics.test_atomic_exch2c                 C   sd   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 )NrB  rC  r   r   r4   r   r=   r>   r/  r   )r   r   r   r   r   r   r   r   r   r   r   rE  r   r   r   test_atomic_exch39  s
    "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 )	NrB  rC  r   r4   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_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r4   r4   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   s   | j tjddd d S N   r   rM  rN  )rP  r   r  r   r   r   r   test_atomic_max_int32S  s    z%TestCudaAtomics.test_atomic_max_int32c                 C   s   | j tjddd d S Nr   rS  rT  )rP  r   r   r   r   r   r   test_atomic_max_uint32V  s    z&TestCudaAtomics.test_atomic_max_uint32c                 C   s   | j tjddd d S rQ  )rP  r   r   r   r   r   r   test_atomic_max_int64Y  s    z%TestCudaAtomics.test_atomic_max_int64c                 C   s   | j tjddd d S rV  )rP  r   r   r   r   r   r   test_atomic_max_uint64\  s    z&TestCudaAtomics.test_atomic_max_uint64c                 C   s   | j tjddd d S rQ  )rP  r   r   r   r   r   r   test_atomic_max_float32_  s    z'TestCudaAtomics.test_atomic_max_float32c                 C   s   | j tjddd d S rQ  )rP  r   r   r   r   r   r   test_atomic_max_doubleb  s    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   rS  rI  r   r*   void(float64[:], float64[:,:]))r   r   r   r   r   r   r   r   !atomic_max_double_normalizedindexrK  r   r   r   rO  r   r   r   r   r   r   &test_atomic_max_double_normalizedindexe  s    
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      r4   r   r*   void(float64[:], float64[:])r   )r   r   r   r   r   r   r   r   atomic_max_double_oneindexrK  r   r   r_  r   r   r   test_atomic_max_double_oneindexo  s    
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 )NrI  r   rS  r   )r   r   r   r   r   r   r   r   
atomic_minminr   r   rL  r   r   r   check_atomic_miny  s    

z TestCudaAtomics.check_atomic_minc                 C   s   | j tjddd d S rQ  )rh  r   r  r   r   r   r   test_atomic_min_int32  s    z%TestCudaAtomics.test_atomic_min_int32c                 C   s   | j tjddd d S rV  )rh  r   r   r   r   r   r   test_atomic_min_uint32  s    z&TestCudaAtomics.test_atomic_min_uint32c                 C   s   | j tjddd d S rQ  )rh  r   r   r   r   r   r   test_atomic_min_int64  s    z%TestCudaAtomics.test_atomic_min_int64c                 C   s   | j tjddd d S rV  )rh  r   r   r   r   r   r   test_atomic_min_uint64  s    z&TestCudaAtomics.test_atomic_min_uint64c                 C   s   | j tjddd d S rQ  )rh  r   r   r   r   r   r   test_atomic_min_float  s    z%TestCudaAtomics.test_atomic_min_floatc                 C   s   | j tjddd d S rQ  )rh  r   r   r   r   r   r   test_atomic_min_double  s    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_normalizedindexrg  r   r   r_  r   r   r   &test_atomic_min_double_normalizedindex  s    
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 ra  )r   r   r   r   r   ro  r   r   atomic_min_double_oneindexrg  r   r   r_  r   r   r   test_atomic_min_double_oneindex  s    
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   rb  r*   r*   r   r*   )r   r   r   r   r   r   r   r   nanr   r   )r   r   r   rO  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   rb  r*   r   rt  )r   r   r   r   r   r   r   r   r   ru  r   r   )r   r   r   r   r   rO  r   r   r   _test_atomic_minmax_nan_val  s    z+TestCudaAtomics._test_atomic_minmax_nan_valc                 C   s   |  t d S r   )rv  rf  r   r   r   r   test_atomic_min_nan_location  s    z,TestCudaAtomics.test_atomic_min_nan_locationc                 C   s   |  t d S r   )rv  rJ  r   r   r   r   test_atomic_max_nan_location  s    z,TestCudaAtomics.test_atomic_max_nan_locationc                 C   s   |  t d S r   )rw  rf  r   r   r   r   test_atomic_min_nan_val  s    z'TestCudaAtomics.test_atomic_min_nan_valc                 C   s   |  t d S r   )rw  rJ  r   r   r   r   test_atomic_max_nan_val  s    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   r4   r   r*   rc  r   )r   r   r   r   r   r   r   r   atomic_max_double_sharedrK  r   r   r   rO  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   ro  r   r   atomic_min_double_sharedrg  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   shuffleZasarrayZ
zeros_liker   r   r   r   r   r   r   r   Zassert_array_equal)r   nfillunfillr   r   r   r   Z	fill_maskZunfill_maskZ
expect_resZ
expect_outr   r   r   r   check_compare_and_swap  s"    



z&TestCudaAtomics.check_compare_and_swapc                 C   s   | j dddtjd d S )NrC  ir  r  r  r   )r  r   r  r   r   r   r   test_atomic_compare_and_swap  s    z,TestCudaAtomics.test_atomic_compare_and_swapc                 C   s   | j dddtjd d S )NrC  ir  r  )r  r   r   r   r   r   r   test_atomic_compare_and_swap2  s    z-TestCudaAtomics.test_atomic_compare_and_swap2c                 C   s@   t jjddt jd}t jjddt jd}| jd||t jd d S NrB  r   r   r*      rC  r  )r   r   r   r   r  r   ZrfillZrunfillr   r   r   test_atomic_compare_and_swap3  s
    
z-TestCudaAtomics.test_atomic_compare_and_swap3c                 C   s@   t jjddt jd}t jjddt jd}| jd||t jd d S r  )r   r   r   r   r  r  r   r   r   test_atomic_compare_and_swap4
  s
    
z-TestCudaAtomics.test_atomic_compare_and_swap4c                 C   sX   t jdt jd}||d< |d | t |rD| t |d  n| |d | d S )Nr  r   r   rt  r*   )r   r   r   isnanr   assertEqualr   r   initialr   r   r   r   _test_atomic_returns_old  s    
z(TestCudaAtomics._test_atomic_returns_oldc                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dd| d< d S r)   )r   r6   r7   r   r   r   r   r     s    z;TestCudaAtomics.test_atomic_add_returns_old.<locals>.kernelr  r   r   r  r   r   r   r   r   test_atomic_add_returns_old  s    
z+TestCudaAtomics.test_atomic_add_returns_oldc                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dd| d< d S r)   r   r6   rK  r  r   r   r   r   &  s    zBTestCudaAtomics.test_atomic_max_returns_no_replace.<locals>.kernelr  r  r  r   r   r   "test_atomic_max_returns_no_replace%  s    
z2TestCudaAtomics.test_atomic_max_returns_no_replacec                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dd| d< d S Nr   r  r*   r  r  r   r   r   r   -  s    zCTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernelr*   r  r  r   r   r   #test_atomic_max_returns_old_replace,  s    
z3TestCudaAtomics.test_atomic_max_returns_old_replacec                 C   s    t jdd }| |tj d S )Nc                 S   s   t j| dd| d< d S r)   r  r  r   r   r   r   4  s    zHTestCudaAtomics.test_atomic_max_returns_old_nan_in_array.<locals>.kernelr   r   r  r   ru  r  r   r   r   (test_atomic_max_returns_old_nan_in_array3  s    
z8TestCudaAtomics.test_atomic_max_returns_old_nan_in_arrayc                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dtj| d< d S r)   )r   r6   rK  r   ru  r  r   r   r   r   ;  s    zCTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernelr  r  r  r   r   r   #test_atomic_max_returns_old_nan_val:  s    
z3TestCudaAtomics.test_atomic_max_returns_old_nan_valc                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dd| d< d S Nr      r*   r   r6   rg  r  r   r   r   r   B  s    zFTestCudaAtomics.test_atomic_min_returns_old_no_replace.<locals>.kernelr  r  r  r   r   r   &test_atomic_min_returns_old_no_replaceA  s    
z6TestCudaAtomics.test_atomic_min_returns_old_no_replacec                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dd| d< d S r  r  r  r   r   r   r   I  s    zCTestCudaAtomics.test_atomic_min_returns_old_replace.<locals>.kernelr  r  r  r   r   r   #test_atomic_min_returns_old_replaceH  s    
z3TestCudaAtomics.test_atomic_min_returns_old_replacec                 C   s    t jdd }| |tj d S )Nc                 S   s   t j| dd| d< d S r  r  r  r   r   r   r   P  s    zHTestCudaAtomics.test_atomic_min_returns_old_nan_in_array.<locals>.kernelr  r  r   r   r   (test_atomic_min_returns_old_nan_in_arrayO  s    
z8TestCudaAtomics.test_atomic_min_returns_old_nan_in_arrayc                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dtj| d< d S r)   )r   r6   rg  r   ru  r  r   r   r   r   W  s    zCTestCudaAtomics.test_atomic_min_returns_old_nan_val.<locals>.kernelr  r  r  r   r   r   #test_atomic_min_returns_old_nan_valV  s    
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 )NrI  r   r*   r  r   )r   r   r   r   r   r   r   r   atomic_nanmaxnanmaxr   r   	r   r   rM  rN  init_valrO  r   r   r   r   r   r   check_atomic_nanmax`  s    

z#TestCudaAtomics.check_atomic_nanmaxc                 C   s   | j tjdddd d S NrR  rS  r   r   rM  rN  r  )r  r   r  r   r   r   r   test_atomic_nanmax_int32i  s    z(TestCudaAtomics.test_atomic_nanmax_int32c                 C   s   | j tjdddd d S Nr   rS  r  )r  r   r   r   r   r   r   test_atomic_nanmax_uint32m  s    z)TestCudaAtomics.test_atomic_nanmax_uint32c                 C   s   | j tjdddd d S r  )r  r   r   r   r   r   r   test_atomic_nanmax_int64q  s    z(TestCudaAtomics.test_atomic_nanmax_int64c                 C   s   | j tjdddd d S r  )r  r   r   r   r   r   r   test_atomic_nanmax_uint64u  s    z)TestCudaAtomics.test_atomic_nanmax_uint64c                 C   s   | j tjddtjd d S NrR  rS  r  )r  r   r   ru  r   r   r   r   test_atomic_nanmax_float32y  s    z*TestCudaAtomics.test_atomic_nanmax_float32c                 C   s   | j tjddtjd d S r  )r  r   r   ru  r   r   r   r   test_atomic_nanmax_double}  s    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   r4   r   r*   r  r   rc  r   )r   r   r   r   r   ru  r   r   r   r   atomic_nanmax_double_sharedr  r   r   r~  r   r   r    test_atomic_nanmax_double_shared  s    
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   rb  r4   r   r*   r  rc  r   )r   r   r   r   r   ru  r   r   r   rd  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 )NrI  r   r*   r  rS  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   s   | j tjdddd d S r  )r  r   r  r   r   r   r   test_atomic_nanmin_int32  s    z(TestCudaAtomics.test_atomic_nanmin_int32c                 C   s   | j tjdddd d S r  )r  r   r   r   r   r   r   test_atomic_nanmin_uint32  s    z)TestCudaAtomics.test_atomic_nanmin_uint32c                 C   s   | j tjdddd d S r  )r  r   r   r   r   r   r   test_atomic_nanmin_int64  s    z(TestCudaAtomics.test_atomic_nanmin_int64c                 C   s   | j tjdddd d S r  )r  r   r   r   r   r   r   test_atomic_nanmin_uint64  s    z)TestCudaAtomics.test_atomic_nanmin_uint64c                 C   s   | j tjddtjd d S r  )r  r   r   ru  r   r   r   r   test_atomic_nanmin_float  s    z(TestCudaAtomics.test_atomic_nanmin_floatc                 C   s   | j tjddtjd d S r  )r  r   r   ru  r   r   r   r   test_atomic_nanmin_double  s    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   ru  r   r   r   r   atomic_nanmin_double_sharedr  r   r   r~  r   r   r    test_atomic_nanmin_double_shared  s    
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   ru  r   r   r   rr  r  r   r   r_  r   r   r   "test_atomic_nanmin_double_oneindex  s    
z2TestCudaAtomics.test_atomic_nanmin_double_oneindexc                 C   sv   t jdt jd}||d< t j|d< |d | t |rb| t |d  | t |d  n| |d | d S )Nr  r   r   r*   rt  )r   r   r   ru  r  ZassertFalser   r  r  r   r   r   _test_atomic_nan_returns_old  s    

z,TestCudaAtomics._test_atomic_nan_returns_oldc                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dd| d< d S r)   r   r6   r  r  r   r   r   r     s    zITestCudaAtomics.test_atomic_nanmax_returns_old_no_replace.<locals>.kernelr  r   r   r  r  r   r   r   )test_atomic_nanmax_returns_old_no_replace  s    
z9TestCudaAtomics.test_atomic_nanmax_returns_old_no_replacec                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dd| d< d S r  r  r  r   r   r   r     s    zFTestCudaAtomics.test_atomic_nanmax_returns_old_replace.<locals>.kernelr*   r  r  r   r   r   &test_atomic_nanmax_returns_old_replace  s    
z6TestCudaAtomics.test_atomic_nanmax_returns_old_replacec                 C   s    t jdd }| |tj d S )Nc                 S   s   t j| dd| d< d S r)   r  r  r   r   r   r     s    zKTestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array.<locals>.kernelr   r   r  r   ru  r  r   r   r   +test_atomic_nanmax_returns_old_nan_in_array  s    
z;TestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_arrayc                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dtj| d< d S r)   )r   r6   r  r   ru  r  r   r   r   r     s    zFTestCudaAtomics.test_atomic_nanmax_returns_old_nan_val.<locals>.kernelr  r  r  r   r   r   &test_atomic_nanmax_returns_old_nan_val  s    
z6TestCudaAtomics.test_atomic_nanmax_returns_old_nan_valc                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dd| d< d S r  r   r6   r  r  r   r   r   r     s    zITestCudaAtomics.test_atomic_nanmin_returns_old_no_replace.<locals>.kernelr  r  r  r   r   r   )test_atomic_nanmin_returns_old_no_replace  s    
z9TestCudaAtomics.test_atomic_nanmin_returns_old_no_replacec                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dd| d< d S r  r  r  r   r   r   r     s    zFTestCudaAtomics.test_atomic_nanmin_returns_old_replace.<locals>.kernelr  r  r  r   r   r   &test_atomic_nanmin_returns_old_replace   s    
z6TestCudaAtomics.test_atomic_nanmin_returns_old_replacec                 C   s    t jdd }| |tj d S )Nc                 S   s   t j| dd| d< d S r  r  r  r   r   r   r     s    zKTestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_array.<locals>.kernelr  r  r   r   r   +test_atomic_nanmin_returns_old_nan_in_array  s    
z;TestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_arrayc                 C   s   t jdd }| |d d S )Nc                 S   s   t j| dtj| d< d S r)   )r   r6   r  r   ru  r  r   r   r   r     s    zFTestCudaAtomics.test_atomic_nanmin_returns_old_nan_val.<locals>.kernelr  r  r  r   r   r   &test_atomic_nanmin_returns_old_nan_val  s    
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   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  r2  r3  r4  r5  r6  r7  r8  r9  r:  r;  r<  r=  r>  r?  r@  rA  rD  rF  rG  rH  rP  rU  rW  rX  rY  rZ  r[  r`  re  rh  ri  rj  rk  rl  rm  rn  rq  rs  rv  rw  rx  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  __classcell__r   r   r   r   r     s  
	
				



	
	


	
	
r   __main__)pZnumpyr   textwrapr   Znumbar   r   r   r   r   Znumba.cuda.testingr   r	   r
   r   Z
numba.corer   r   r   r   r   r'   r(   r/   r1   r2   r9   r:   r@   rA   rB   rG   rI   rL   rN   rO   rS   rT   rW   rX   rY   r[   r\   r^   r_   r`   rb   rc   rd   re   rf   rg   rh   ri   rk   rl   rm   rn   rq   rs   rt   rw   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   rJ  r^  rd  r}  rf  rp  rr  r  r  Z$atomic_nanmax_double_normalizedindexZatomic_nanmax_double_oneindexr  r  Z$atomic_nanmin_double_normalizedindexZatomic_nanmin_double_oneindexr  r   r   r  mainr   r   r   r   <module>   s   













	

	%          S
