U
    hd                     @   s  d dl Zd dlmZ d dlmZmZmZmZm	Z	 d dl
mZmZmZ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 Zldd Zmdd ZnG dd deZoepdkr~eq  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   ]/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/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   fnsldr   r   r   gen_atomic_extreme_funcs  s     r   zcuda.atomic.maxzcuda.atomic.minzcuda.atomic.nanmaxzcuda.atomic.nanminc                 C   s8   t d}|| jk r4t j| |d  ||| ||< d S Nr*   )r   gridsizer6   Zcompare_and_swapresoldr   Zfill_valgidr   r   r   atomic_compare_and_swap  s    

r   c                 C   s2   t d}|| jk r.t j| |||| ||< d S r   )r   r   r   r6   casr   r   r   r   atomic_cas_1dim  s    

r   c                 C   sL   t d}|d | jd k rH|d | jd k rHt j| |||| ||< d S )N   r   r*   )r   r   r0   r6   r   r   r   r   r   atomic_cas_2dim  s    
$r   c                       s  e Zd Z fddZdd Zdd Zdd Zd	d
 Zdd Zdd Z	d)ddZ
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Ä Zdd*ddƄ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 Zdd  Zdd Zdd Zdd Zdd Zd	d
 Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd  Zd!d" Zd#d$ Zd%d& Zd'd( Z  ZS (+  TestCudaAtomicsc                    s   t    tjd d S )Nr   )supersetUpnprandomseedself	__class__r   r   r     s    
zTestCudaAtomics.setUpc                 C   s   t jjddddt j}| }| }tdt}|d | tdt	}|d | t j
dt jd}t|jD ]}|||   d7  < 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   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_double<  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_2Q  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_3a  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_globalj  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_2  s    "z/TestCudaAtomics.test_atomic_add_double_global_2c                 C   sh   t jjddddt jdd}| }tdt	}|d | t j
||d  | j|d	d
 d S r   )r   r   r   r   r   r   r   r   r   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_or6  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_or2C  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_or3K  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_globalS  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_2a  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_xori  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_xor2v  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_xor3~  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   r4  r   r   r   test_atomic_inc2_64  s    z#TestCudaAtomics.test_atomic_inc2_64c                 C   s,   |  tj\}}d}| |||ddt d S r3  )r#  r   r   r,  r   r4  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 r1  )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 r3  )r#  r   r   r,  r   r4  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 r6  )r#  r   r   r,  r   r4  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   r>  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 r1  )r"  r   r   r>  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 r3  )r#  r   r   r@  r   r4  r   r   r   test_atomic_dec2_32  s    z#TestCudaAtomics.test_atomic_dec2_32c                 C   s,   |  tj\}}d}| |||ddt d S r6  )r#  r   r   r@  r   r4  r   r   r   test_atomic_dec2_64  s    z#TestCudaAtomics.test_atomic_dec2_64c                 C   s,   |  tj\}}d}| |||ddt d S r3  )r#  r   r   r@  r   r4  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   r?  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 r1  )r"  r   r   r?  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 r3  )r#  r   r   r@  r   r4  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 r6  )r#  r   r   r@  r   r4  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_exch2  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 )NrJ  rK  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_exch2<  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 )NrJ  rK  r   r   r4   r   r=   r>   r7  r   )r   r   r   r   r   r   r   r   r   r   r   rM  r   r   r   test_atomic_exch3D  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 )	NrJ  rK  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_globalL  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_maxV  s    

z TestCudaAtomics.check_atomic_maxc                 C   s   | j tjddd d S N   r   rU  rV  )rX  r   r  r   r   r   r   test_atomic_max_int32^  s    z%TestCudaAtomics.test_atomic_max_int32c                 C   s   | j tjddd d S Nr   r[  r\  )rX  r   r   r   r   r   r   test_atomic_max_uint32a  s    z&TestCudaAtomics.test_atomic_max_uint32c                 C   s   | j tjddd d S rY  )rX  r   r   r   r   r   r   test_atomic_max_int64d  s    z%TestCudaAtomics.test_atomic_max_int64c                 C   s   | j tjddd d S r^  )rX  r   r   r   r   r   r   test_atomic_max_uint64g  s    z&TestCudaAtomics.test_atomic_max_uint64c                 C   s   | j tjddd d S rY  )rX  r   r   r   r   r   r   test_atomic_max_float32j  s    z'TestCudaAtomics.test_atomic_max_float32c                 C   s   | j tjddd d S rY  )rX  r   r   r   r   r   r   test_atomic_max_doublem  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   r[  rQ  r   r*   void(float64[:], float64[:,:]))r   r   r   r   r   r   r   r   !atomic_max_double_normalizedindexrS  r   r   r   rW  r   r   r   r   r   r   &test_atomic_max_double_normalizedindexp  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_oneindexrS  r   r   rg  r   r   r   test_atomic_max_double_oneindexz  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 )NrQ  r   r[  r   )r   r   r   r   r   r   r   r   
atomic_minminr   r   rT  r   r   r   check_atomic_min  s    

z TestCudaAtomics.check_atomic_minc                 C   s   | j tjddd d S rY  )rp  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 r^  )rp  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 rY  )rp  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 r^  )rp  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 rY  )rp  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 rY  )rp  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 rd  )r   r   r   r   r   onesr   r   !atomic_min_double_normalizedindexro  r   r   rg  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 ri  )r   r   r   r   r   rw  r   r   atomic_min_double_oneindexro  r   r   rg  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 )Nre  r   rj  r*   r*   r   r*   )r   r   r   r   r   r   r   r   nanr   r   )r   r   r   rW  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 )Nre  r   rj  r*   r   r|  )r   r   r   r   r   r   r   r   r   r}  r   r   )r   r   r   r   r   rW  r   r   r   _test_atomic_minmax_nan_val  s    z+TestCudaAtomics._test_atomic_minmax_nan_valc                 C   s   |  t d S r   )r~  rn  r   r   r   r   test_atomic_min_nan_location  s    z,TestCudaAtomics.test_atomic_min_nan_locationc                 C   s   |  t d S r   )r~  rR  r   r   r   r   test_atomic_max_nan_location  s    z,TestCudaAtomics.test_atomic_max_nan_locationc                 C   s   |  t d S r   )r  rn  r   r   r   r   test_atomic_min_nan_val  s    z'TestCudaAtomics.test_atomic_min_nan_valc                 C   s   |  t d S r   )r  rR  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*   rk  r   )r   r   r   r   r   r   r   r   atomic_max_double_sharedrS  r   r   r   rW  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   rw  r   r   atomic_min_double_sharedro  r   r   r  r   r   r   test_atomic_min_double_shared  s    
z-TestCudaAtomics.test_atomic_min_double_sharedr*   c                 C   s   |g|d  |g|d   }t j| t j||d}|dkrDd|_t |}t jjdd|jd|j}	||k}
||k}t |}|	|
 ||
< |||< |	 }t
|}|dkr|d |||	| n|d |||	| t j|| t j|| d S )	Nr   r   )
   r*   r  r   r  r  )r  r  )r   r   shuffleZasarrayr0   Z
zeros_liker   r   r   r   r   r   r   Zassert_array_equal)r   nfillunfillr   cas_funcndimr   outr   Z	fill_maskZunfill_maskZ
expect_resZ
expect_outr   r   r   r   	check_cas  s&    


zTestCudaAtomics.check_casc                 C   s   | j dddtjtd d S NrK  r  r  r  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td d S NrK  r  r  )r  r   r   r   r   r   r   r   test_atomic_compare_and_swap2  s    z-TestCudaAtomics.test_atomic_compare_and_swap2c                 C   sB   t jjddt jd}t jjddt jd}| jd||t jtd d S NrJ  r  r   r*      rK  r  )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   sB   t jjddt jd}t jjddt jd}| jd||t jtd d S r  )r   r   r   r   r  r   r  r   r   r   test_atomic_compare_and_swap4  s
    z-TestCudaAtomics.test_atomic_compare_and_swap4c                 C   s   | j dddtjtd d S r  )r  r   r  r   r   r   r   r   test_atomic_cas_1dim   s    z$TestCudaAtomics.test_atomic_cas_1dimc                 C   s   | j dddtjtdd d S )NrK  r  r  r   r  r  r  r   r  r  )r  r   r  r   r   r   r   r   test_atomic_cas_2dim$  s     z$TestCudaAtomics.test_atomic_cas_2dimc                 C   s   | j dddtjtd d S r  )r  r   r   r   r   r   r   r   test_atomic_cas2_1dim(  s    z%TestCudaAtomics.test_atomic_cas2_1dimc                 C   s   | j dddtjtdd d S )NrK  r  r  r   r  )r  r   r   r   r   r   r   r   test_atomic_cas2_2dim,  s     z%TestCudaAtomics.test_atomic_cas2_2dimc                 C   sB   t jjddt jd}t jjddt jd}| jd||t jtd d S r  )r   r   r   r   r  r   r  r   r   r   test_atomic_cas3_1dim0  s
    z%TestCudaAtomics.test_atomic_cas3_1dimc                 C   sD   t jjddt jd}t jjddt jd}| jd||t jtdd d S 	NrJ  r  r   r*   r  rK  r   r  )r   r   r   r   r  r   r  r   r   r   test_atomic_cas3_2dim6  s     z%TestCudaAtomics.test_atomic_cas3_2dimc                 C   sB   t jjddt jd}t jjddt jd}| jd||t jtd d S r  )r   r   r   r   r  r   r  r   r   r   test_atomic_cas4_1dim<  s
    z%TestCudaAtomics.test_atomic_cas4_1dimc                 C   sD   t jjddt jd}t jjddt jd}| jd||t jtdd d S r  )r   r   r   r   r  r   r  r   r   r   test_atomic_cas4_2dimB  s     z%TestCudaAtomics.test_atomic_cas4_2dimc                 C   sX   t jdt jd}||d< |d | t |rD| t |d  n| |d | d S )Nr   r   r   r|  r*   )r   r   r   isnanr   assertEqualr   r   initialr   r   r   r   _test_atomic_returns_oldM  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   W  s    z;TestCudaAtomics.test_atomic_add_returns_old.<locals>.kernelr  r   r   r  r   r   r   r   r   test_atomic_add_returns_oldV  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   rS  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   e  s    zCTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernelr*   r  r  r   r   r   #test_atomic_max_returns_old_replaced  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   l  s    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_arrayk  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   rS  r   r}  r  r   r   r   r   s  s    zCTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernelr  r  r  r   r   r   #test_atomic_max_returns_old_nan_valr  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   ro  r  r   r   r   r   z  s    zFTestCudaAtomics.test_atomic_min_returns_old_no_replace.<locals>.kernelr  r  r  r   r   r   &test_atomic_min_returns_old_no_replacey  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     s    zCTestCudaAtomics.test_atomic_min_returns_old_replace.<locals>.kernelr  r  r  r   r   r   #test_atomic_min_returns_old_replace  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     s    zHTestCudaAtomics.test_atomic_min_returns_old_nan_in_array.<locals>.kernelr  r  r   r   r   (test_atomic_min_returns_old_nan_in_array  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   ro  r   r}  r  r   r   r   r     s    zCTestCudaAtomics.test_atomic_min_returns_old_nan_val.<locals>.kernelr  r  r  r   r   r   #test_atomic_min_returns_old_nan_val  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 )NrQ  r   r*   r   r   )r   r   r   r   r   r   r   r   atomic_nanmaxnanmaxr   r   	r   r   rU  rV  init_valrW  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 NrZ  r[  r   r   rU  rV  r  )r  r   r  r   r   r   r   test_atomic_nanmax_int32  s    z(TestCudaAtomics.test_atomic_nanmax_int32c                 C   s   | j tjdddd d S Nr   r[  r  )r  r   r   r   r   r   r   test_atomic_nanmax_uint32  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_int64  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_uint64  s    z)TestCudaAtomics.test_atomic_nanmax_uint64c                 C   s   | j tjddtjd d S NrZ  r[  r  )r  r   r   r}  r   r   r   r   test_atomic_nanmax_float32  s    z*TestCudaAtomics.test_atomic_nanmax_float32c                 C   s   | j tjddtjd d S r  )r  r   r   r}  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   rk  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  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   rj  r4   r   r*   r   rk  r   )r   r   r   r   r   r}  r   r   r   rl  r  r   r   rg  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 )NrQ  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   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   r}  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   r}  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   r}  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   r}  r   r   r   rz  r  r   r   rg  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*   r|  )r   r   r   r}  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   r}  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   r}  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   2  s    zITestCudaAtomics.test_atomic_nanmin_returns_old_no_replace.<locals>.kernelr  r  r  r   r   r   )test_atomic_nanmin_returns_old_no_replace1  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   9  s    zFTestCudaAtomics.test_atomic_nanmin_returns_old_replace.<locals>.kernelr  r  r  r   r   r   &test_atomic_nanmin_returns_old_replace8  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   r}  r  r   r   r   r   G  s    zFTestCudaAtomics.test_atomic_nanmin_returns_old_nan_val.<locals>.kernelr  r  r  r   r   r   &test_atomic_nanmin_returns_old_nan_valF  s    
z6TestCudaAtomics.test_atomic_nanmin_returns_old_nan_val)T)r*   )__name__
__module____qualname__r   r   r   r   r   r   r   r   r
   r   r   r   r   r   r   r   r   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  r2  r5  r8  r9  r:  r;  r<  r=  r>  r?  r@  rA  rB  rC  rD  rE  rF  rG  rH  rI  rL  rN  rO  rP  rX  r]  r_  r`  ra  rb  rc  rh  rm  rp  rq  rr  rs  rt  ru  rv  ry  r{  r~  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  __classcell__r   r   r   r   r     s*  
	
				



	
	


	
	

r   __main__)r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   rR  rf  rl  r  rn  rx  rz  r  r  Z$atomic_nanmax_double_normalizedindexZatomic_nanmax_double_oneindexr  r  Z$atomic_nanmin_double_normalizedindexZatomic_nanmin_double_oneindexr  r   r   r   r   r  mainr   r   r   r   <module>   s   













	

	%           
