U
    -eQ                     @   s~  d dl Zd dlmZ d dlmZmZmZmZm	Z	 d dl
mZmZmZ d dlmZ ejdddd	 Zejddd
d Zejdddd Zejdddd Zejdddd Zejdddd Zejdddd Zejdddd Zdd Zdd Zdd Zdd Zd d! Zd"d# Zd$d% Zd&d' Z d(d) Z!d*d+ Z"d,d- Z#d.d/ Z$d0d1 Z%d2d3 Z&d4d5 Z'd6d7 Z(d8d9 Z)d:d; Z*d<d= Z+d>d? Z,d@dA Z-dBdC Z.dDdE Z/dFdG Z0dHdI Z1dJdK Z2dLdM Z3dNdO Z4dPdQ Z5dRdS Z6dTdU Z7dVdW Z8dXdY Z9dZd[ Z:d\d] Z;d^d_ Z<d`da Z=dbdc Z>ddde Z?dfdg Z@dhdi ZAdjdk ZBdldm ZCdndo ZDdpdq ZEdrds ZFdtdu ZGdvdw ZHdxdy ZIdzd{ ZJd|d} ZKd~d ZLdd ZMdd ZNdd ZOdd ZPdd ZQdd ZRdd ZSdd ZTdd ZUdd ZVdd ZWdd ZXdd ZYdd ZZeZd\Z[Z\Z]Z^eZd\Z_Z`ZaZbeZd\ZcZdZeZfeZd\ZgZhZiZjdd Zkdd Zldd ZmG dd deZneodkrzep  dS )    N)dedent)cudauint32uint64float32float64)unittestCUDATestCasecc_X_or_above)configT)Zdevicec                 C   s   t | S N)r   num r   e/var/www/html/Darija-Ai-Train/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_shared5   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_globalE   s    $r0   c                 C   s4   t jj}t|| | }|r$|| }|| || d S r   )r   r   r   r   )r   r   r   r   r    r"   r#   r%   r   r   r   atomic_binary_1dim_globalO   s
    r1   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_addY   s       r8   c              
   C   s    t | | dtdtjjtdd	 d S )Nr)   r3   r   Tr4   r7   r   r   r   atomic_add_wrap^   s       r9   c                 C   s   t | dtdtjjtd d S Nr)         Fr.   r   r   r5   r6   r   r7   r   r   r   atomic_add2c   s
    
  r?   c                 C   s   t | dtdtjjtd d S )Nr)   r;   Tr>   r7   r   r   r   atomic_add2_wraph   s
    
  r@   c                 C   s   t | dtdtjjtd d S r:   )r.   r   r   r5   r6   r   r7   r   r   r   atomic_add3m   s
    
  rA   c              
   C   s    t | | dtdtjjtdd	 d S N      ?r3           Fr&   r   r   r5   r6   r   r7   r   r   r   atomic_add_floatr   s       rF   c              
   C   s    t | | dtdtjjtdd	 d S NrC   r3   rD   TrE   r7   r   r   r   atomic_add_float_wrapw   s       rH   c                 C   s   t | dtdtjjtd d S NrC   r;   Fr.   r   r   r5   r6   r   r7   r   r   r   atomic_add_float_2|   s
    
  rK   c                 C   s   t | dtdtjjtd d S NrC   r;   TrJ   r7   r   r   r   atomic_add_float_2_wrap   s
    
  rM   c                 C   s   t | dtdtjjtd d S rI   )r.   r   r   r5   r6   r   r7   r   r   r   atomic_add_float_3   s
    
  rN   c                 C   s   t || ddtjjd d S Nr3   rC   Fr1   r   r5   r6   r   r   r   r   r   atomic_add_double_global   s    rR   c                 C   s   t || ddtjjd d S )Nr3   rC   TrP   rQ   r   r   r   atomic_add_double_global_wrap   s    rS   c                 C   s   t | dtjjtd d S Nr)   Fr0   r   r5   r6   r   r7   r   r   r   atomic_add_double_global_2   s    rV   c                 C   s   t | dtjjtd d S )Nr)   TrU   r7   r   r   r   atomic_add_double_global_2_wrap   s    rW   c                 C   s   t | dtjjtd d S rT   )r0   r   r5   r6   r   r7   r   r   r   atomic_add_double_global_3   s    rX   c              
   C   s    t || dtdtjjtdd	 d S rB   r&   r   r   r5   r6   r   rQ   r   r   r   atomic_add_double   s       rZ   c              
   C   s    t || dtdtjjtdd	 d S rG   rY   rQ   r   r   r   atomic_add_double_wrap   s       r[   c                 C   s   t | dtdtjjtd d S rI   r.   r   r   r5   r6   r   r7   r   r   r   atomic_add_double_2   s
    
  r]   c                 C   s   t | dtdtjjtd d S rL   r\   r7   r   r   r   atomic_add_double_2_wrap   s
    
  r^   c                 C   s   t | dtdtjjtd d S rI   )r.   r   r   r5   r6   r   r7   r   r   r   atomic_add_double_3   s
    
  r_   c              
   C   s    t | | dtdtjjtdd	 d S r2   )r&   r   r   r5   subr   r7   r   r   r   
atomic_sub   s       ra   c                 C   s   t | dtdtjjtd d S r:   )r.   r   r   r5   r`   r   r7   r   r   r   atomic_sub2   s
    
  rb   c                 C   s   t | dtdtjjtd d S r:   )r.   r   r   r5   r`   r   r7   r   r   r   atomic_sub3   s
    
  rc   c              
   C   s    t | | dtdtjjtdd	 d S rB   )r&   r   r   r5   r`   r   r7   r   r   r   atomic_sub_float   s       rd   c                 C   s   t | dtdtjjtd d S rI   )r.   r   r   r5   r`   r   r7   r   r   r   atomic_sub_float_2   s
    
  re   c                 C   s   t | dtdtjjtd d S rI   )r.   r   r   r5   r`   r   r7   r   r   r   atomic_sub_float_3   s
    
  rf   c              
   C   s    t || dtdtjjtdd	 d S rB   )r&   r   r   r5   r`   r   rQ   r   r   r   atomic_sub_double   s       rg   c                 C   s   t | dtdtjjtd d S rI   )r.   r   r   r5   r`   r   r7   r   r   r   atomic_sub_double_2   s
    
  rh   c                 C   s   t | dtdtjjtd d S rI   r.   r   r   r5   r`   r   r7   r   r   r   atomic_sub_double_3   s
    
  rj   c                 C   s   t || ddtjjd d S rO   )r1   r   r5   r`   rQ   r   r   r   atomic_sub_double_global   s    rk   c                 C   s   t | dtjjtd d S )NrC   F)r0   r   r5   r`   r   r7   r   r   r   atomic_sub_double_global_2   s    rl   c                 C   s   t | dtdtjjtd d S rI   ri   r7   r   r   r   atomic_sub_double_global_3   s
    
  rm   c              
   C   s    t | | |tdtjjtdd	 d S )Nr3   r)   F)r&   r   r   r5   and_r   r   r   r   r   r   
atomic_and   s       rp   c                 C   s   t | |tdtjjtd d S Nr;   F)r.   r   r   r5   rn   r   ro   r   r   r   atomic_and2   s
    
  rr   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   rn   r   ro   r   r   r   atomic_and3   s
    
  rs   c                 C   s   t || d|tjjd d S Nr3   F)r1   r   r5   rn   r   r   r   r   r   r   atomic_and_global  s    rv   c                 C   s   t | |tjjtd d S NF)r0   r   r5   rn   r   ro   r   r   r   atomic_and_global_2  s     rx   c              
   C   s    t | | |tdtjjtdd	 d S Nr3   r   F)r&   r   r   r5   or_r   ro   r   r   r   	atomic_or  s       r{   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   rz   r   ro   r   r   r   
atomic_or2  s
    
  r|   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   rz   r   ro   r   r   r   
atomic_or3  s
    
  r}   c                 C   s   t || d|tjjd d S rt   )r1   r   r5   rz   ru   r   r   r   atomic_or_global  s    r~   c                 C   s   t | |tjjtd d S rw   )r0   r   r5   rz   r   ro   r   r   r   atomic_or_global_2  s     r   c              
   C   s    t | | |tdtjjtdd	 d S ry   )r&   r   r   r5   xorr   ro   r   r   r   
atomic_xor$  s       r   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   r   r   ro   r   r   r   atomic_xor2)  s
    
  r   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   r   r   ro   r   r   r   atomic_xor3.  s
    
  r   c                 C   s   t || d|tjjd d S rt   )r1   r   r5   r   ru   r   r   r   atomic_xor_global3  s    r   c                 C   s   t | |tjjtd d S rw   )r0   r   r5   r   r   ro   r   r   r   atomic_xor_global_27  s     r   c                 C   s   t | ||tdtjjt d S Nr3   )r'   r   r   r5   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   r5   r   r   r   r   r   r   atomic_inc64A  s     r   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   r   r   ro   r   r   r   atomic_inc2_32F  s
    
  r   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   r   r   ro   r   r   r   atomic_inc2_64K  s
    
  r   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   r   r   ro   r   r   r   atomic_inc3P  s
    
  r   c                 C   s   t || d|tjjd d S rt   )r1   r   r5   r   ru   r   r   r   atomic_inc_globalU  s    r   c                 C   s   t | |tjjtd d S rw   )r0   r   r5   r   r   ro   r   r   r   atomic_inc_global_2Y  s     r   c                 C   s   t | ||tdtjjt d S r   )r'   r   r   r5   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   r5   r   r   r   r   r   r   atomic_dec64c  s     r   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   r   r   ro   r   r   r   atomic_dec2_32h  s
    
  r   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   r   r   ro   r   r   r   atomic_dec2_64m  s
    
  r   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   r   r   ro   r   r   r   atomic_dec3r  s
    
  r   c                 C   s   t || d|tjjd d S rt   )r1   r   r5   r   ru   r   r   r   atomic_dec_globalw  s    r   c                 C   s   t | |tjjtd d S rw   )r0   r   r5   r   r   ro   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   r5   exchr   r   r   r   r   atomic_exch  s     r   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   r   r   ro   r   r   r   atomic_exch2  s
    
  r   c                 C   s   t | |tdtjjtd d S rq   )r.   r   r   r5   r   r   ro   r   r   r   atomic_exch3  s
    
  r   c                 C   s   t || d|tjjd d S rt   )r1   r   r5   r   ru   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   r5   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sizer5   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   r5   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   r/   r5   r   r   r   r   r   atomic_cas_2dim  s    
$r   c                       s  e Zd Z fddZdd Zdd Zdd Zd	d
 Zdd Zdd Z	d)ddZ
dd Zdd Zdd Zdd Zdd Zdd Zdd Zd d! Zd"d# Zd$d% Zd&d' Zd(d) Zd*d+ Zd,d- Zd.d/ Zd0d1 Zd2d3 Zd4d5 Zd6d7 Zd8d9 Zd:d; Zd<d= Z d>d? Z!d@dA Z"dBdC Z#dDdE Z$dFdG Z%dHdI Z&dJdK Z'dLdM Z(dNdO Z)dPdQ Z*dRdS Z+dTdU Z,dVdW Z-dXdY Z.dZd[ Z/d\d] Z0d^d_ Z1d`da Z2dbdc Z3ddde Z4dfdg Z5dhdi Z6djdk Z7dldm Z8dndo Z9dpdq Z:drds Z;dtdu Z<dvdw Z=dxdy Z>dzd{ Z?d|d} Z@d~d ZAdd ZBdd ZCdd ZDdd ZEdd ZFdd ZGdd ZHdd ZIdd ZJdd ZKdd ZLdd ZMdd ZNdd ZOdd ZPdd ZQdd ZRdd ZSdd ZTdd ZUdd ZVdd ZWdd ZXdd ZYdd ZZdd Z[dd Z\dd Z]dd Z^dd Z_dd Z`dd Zadd ZbddÄ Zcd*ddƄZdddȄ Zeddʄ Zfdd̄ Zgdd΄ ZhddЄ Zidd҄ ZjddԄ Zkddք Zldd؄ Zmddڄ Zndd܄ Zoddބ Zpdd Zqdd Zrdd Zsdd Ztdd Zudd Zvdd Zwdd Zxdd Zydd Zzdd Z{dd Z|dd Z}dd Z~dd Zdd Zdd  Zdd Zdd Zdd Zdd Zd	d
 Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd  Zd!d" Zd#d$ Zd%d& Zd'd( Z  ZS (+  TestCudaAtomicsc                    s   t    tjd d S )Nr   )supersetUpnprandomseedself	__class__r   r   r     s    
zTestCudaAtomics.setUpc                 C   s   t jjddddt j}| }| }tdt}|d | tdt	}|d | t j
dt jd}t|jD ]}|||   d7  < qx| t ||k | t ||k d S Nr   r3   r   zvoid(uint32[:])r)   r3   dtyper)   )r   r   randintastyper   copyr   jitr8   r9   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   r3   r   r<   r=   zvoid(uint32[:,:])r)   r;   r)   )r   r   r   r   r   reshaper   r   r   r?   r@   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   rA   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   r3   r   zvoid(float32[:])r   r   rC   )r   r   r   r   r   r   intpr   r   rF   rH   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   r3   r   r<   r=   zvoid(float32[:,:])r   r)   )r   r   r   r   r   r   r   r   r   rK   rM   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   rN   r   r   r   r   r   r   test_atomic_add_float_3"  s
    "z'TestCudaAtomics.test_atomic_add_float_3Tc                 C   s   t jr
d S tt|  }tddrbtj	 dkr<d}nd}|rN| d}| 
| d| n|rt| 
d| n| 
d	| d S )
N   r   )   r)   redZatomz.sharedz.add.f64zatom.shared.cas.b64zatom.cas.b64)r   ZENABLE_CUDASIMnextiterZinspect_asmvaluesr
   r   Zruntimeget_versionZassertIn)r   kernelr   asminstr   r   r   assertCorrectFloat64Atomics*  s    

z+TestCudaAtomics.assertCorrectFloat64Atomicsc                 C   s   t jjdddt jd}t dt j}| }tdt	}|d || tdt
}|d || t jdt jd}t|jD ]}|||   d7  < q~t j|| t j|| | | | | d S Nr   r3   r   r   void(int64[:], float64[:])r   r   rC   )r   r   r   int64r   r   r   r   r   rZ   r[   r   r   r   testingassert_equalr   )r   r   r   r   cuda_fnZwrap_fnr   r   r   r   r   test_atomic_add_doubleC  s    
z&TestCudaAtomics.test_atomic_add_doublec                 C   s   t jjddddt jdd}| }| }tdt	}|d | tdt
}|d | t j||d  t j||d  | | | | d S 	Nr   r3   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_2W  s    "
z(TestCudaAtomics.test_atomic_add_double_2c                 C   sd   t jjddddt jdd}| }tdt	}|d | t j
||d  | | d S r   )r   r   r   r   r   r   r   r   r   r_   r   r   r   r   r   r   	cuda_funcr   r   r   test_atomic_add_double_3g  s    "z(TestCudaAtomics.test_atomic_add_double_3c           	      C   s   t jjdddt jd}t dt j}| }d}t|t	}t|t
}|d || |d || t jdt jd}t|jD ]}|||   d7  < qt j|| t j|| | j|dd	 | j|dd	 d S )
Nr   r3   r   r   r   r   rC   Fr   )r   r   r   r   r   r   r   r   r   rR   rS   r   r   r   r   r   r   )	r   r   r   r   sigr   wrap_cuda_funcr   r   r   r   r   test_atomic_add_double_globalp  s    z-TestCudaAtomics.test_atomic_add_double_globalc                 C   s   t jjddddt jdd}| }| }d}t|t	}t|t
}|d | |d | t j||d  t j||d  | j|d	d
 | j|d	d
 d S Nr   r3   r   r<   r=   r   r   r)   Fr   )r   r   r   r   r   r   r   r   r   rV   rW   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   rX   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   ra   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   rb   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   rc   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   rd   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   re   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   rf   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   rg   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   rh   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   rj   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   rk   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   rl   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   rm   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   r3   r   void(uint32[:], uint32)r   )r   r   r   r   r   r   r   r   rp   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   r3   r   r<   r=   void(uint32[:,:], uint32)r   )r   r   r   r   r   r   r   r   r   rr   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   rs   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   r3   r   zvoid(int32[:], int32[:], int32)r   )r   r   r   int32r   r   rv   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   rx   r   r   r   r	  r   r   r   r   r   r   test_atomic_and_global_23  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   r3   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_or2H  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_or3P  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_globalX  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_2f  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_xorn  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_xor2{  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 )Nr3   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 )Nr3   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)   r3   )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)   r3   )r%  r   r   r,  r   r2  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   r7  r   r   r   test_atomic_inc2_64  s    z#TestCudaAtomics.test_atomic_inc2_64c                 C   s,   |  tj\}}d}| |||ddt d S r6  )r&  r   r   r/  r   r7  r   r   r   test_atomic_inc3  s    z TestCudaAtomics.test_atomic_inc3c              	   C   s2   | j tjd\}}}d}| ||||ddt d S r0  )r%  r   r   r-  r   r2  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 r4  )r%  r   r   r-  r   r2  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 r6  )r&  r   r   r/  r   r7  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 r9  )r&  r   r   r/  r   r7  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 r0  )r%  r   r   rA  r   r2  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 r4  )r%  r   r   rA  r   r2  r   r   r   test_atomic_dec_64  s    z"TestCudaAtomics.test_atomic_dec_64c                 C   s,   |  tj\}}d}| |||ddt d S r6  )r&  r   r   rC  r   r7  r   r   r   test_atomic_dec2_32  s    z#TestCudaAtomics.test_atomic_dec2_32c                 C   s,   |  tj\}}d}| |||ddt d S r9  )r&  r   r   rC  r   r7  r   r   r   test_atomic_dec2_64  s    z#TestCudaAtomics.test_atomic_dec2_64c                 C   s,   |  tj\}}d}| |||ddt d S r6  )r&  r   r   rC  r   r7  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 r0  )r%  r   r   rB  r   r2  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 r4  )r%  r   r   rB  r   r2  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 r6  )r&  r   r   rC  r   r7  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 r9  )r&  r   r   rC  r   r7  r   r   r   test_atomic_dec_global2_642  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   r3   r   r1  r   )r   r   r   r   r   r"  r   r   r   r   r   )r   r	  r   r   r   r   r   r   test_atomic_exch7  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 )NrM  rN  r   r   r3   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_exch2A  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 )NrM  rN  r   r   r3   r   r<   r=   r:  r   )r   r   r   r   r   r   r   r   r   r   r   rP  r   r   r   test_atomic_exch3I  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 )	NrM  rN  r   r3   r   r   r1  r   )
r   r   r   r   r"  r   r   r   r   r   )r   r	  r   r   r   r   r   r   r   test_atomic_exch_globalQ  s    z'TestCudaAtomics.test_atomic_exch_globalc                 C   s\   t jj||dd|}t jd|jd}tt}|d || t 	|}t j
|| d S )Nr3   r3   r   r)   r   )r   r   r   r   r   r   r   r   
atomic_maxmaxr   r   r   r   lohivalsr   r   r   r   r   r   check_atomic_max[  s    

z TestCudaAtomics.check_atomic_maxc                 C   s   | j tjddd d S N   r   rX  rY  )r[  r   r  r   r   r   r   test_atomic_max_int32c  s    z%TestCudaAtomics.test_atomic_max_int32c                 C   s   | j tjddd d S Nr   r^  r_  )r[  r   r   r   r   r   r   test_atomic_max_uint32f  s    z&TestCudaAtomics.test_atomic_max_uint32c                 C   s   | j tjddd d S r\  )r[  r   r   r   r   r   r   test_atomic_max_int64i  s    z%TestCudaAtomics.test_atomic_max_int64c                 C   s   | j tjddd d S ra  )r[  r   r   r   r   r   r   test_atomic_max_uint64l  s    z&TestCudaAtomics.test_atomic_max_uint64c                 C   s   | j tjddd d S r\  )r[  r   r   r   r   r   r   test_atomic_max_float32o  s    z'TestCudaAtomics.test_atomic_max_float32c                 C   s   | j tjddd d S r\  )r[  r   r   r   r   r   r   test_atomic_max_doubler  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^  rT  r   r)   void(float64[:], float64[:,:]))r   r   r   r   r   r   r   r   !atomic_max_double_normalizedindexrV  r   r   r   rZ  r   r   r   r   r   r   &test_atomic_max_double_normalizedindexu  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      r3   r   r)   void(float64[:], float64[:])r   )r   r   r   r   r   r   r   r   atomic_max_double_oneindexrV  r   r   rj  r   r   r   test_atomic_max_double_oneindex  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 )NrT  r   r^  r   )r   r   r   r   r   r   r   r   
atomic_minminr   r   rW  r   r   r   check_atomic_min  s    

z TestCudaAtomics.check_atomic_minc                 C   s   | j tjddd d S r\  )rs  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 ra  )rs  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 r\  )rs  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 ra  )rs  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 r\  )rs  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 r\  )rs  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 rg  )r   r   r   r   r   onesr   r   !atomic_min_double_normalizedindexrr  r   r   rj  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 rl  )r   r   r   r   r   rz  r   r   atomic_min_double_oneindexrr  r   r   rj  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 )Nrh  r   rm  r)   r)   r   r)   )r   r   r   r   r   r   r   r   nanr   r   )r   r   r   rZ  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 )Nrh  r   rm  r)   r   r  )r   r   r   r   r   r   r   r   r   r  r   r   )r   r   r   r   r   rZ  r   r   r   _test_atomic_minmax_nan_val  s    z+TestCudaAtomics._test_atomic_minmax_nan_valc                 C   s   |  t d S r   )r  rq  r   r   r   r   test_atomic_min_nan_location  s    z,TestCudaAtomics.test_atomic_min_nan_locationc                 C   s   |  t d S r   )r  rU  r   r   r   r   test_atomic_max_nan_location  s    z,TestCudaAtomics.test_atomic_max_nan_locationc                 C   s   |  t d S r   )r  rq  r   r   r   r   test_atomic_min_nan_val  s    z'TestCudaAtomics.test_atomic_min_nan_valc                 C   s   |  t d S r   )r  rU  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   r3   r   r)   rn  r   )r   r   r   r   r   r   r   r   atomic_max_double_sharedrV  r   r   r   rZ  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   rz  r   r   atomic_min_double_sharedrr  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asarrayr/   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 NrN  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 NrN  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 NrM  r  r   r)      rN  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 )NrN  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 )NrN  r  r  r   r  )r  r   r   r   r   r   r   r   test_atomic_cas2_2dim1  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_1dim5  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 	NrM  r  r   r)   r  rN  r   r  )r   r   r   r   r  r   r  r   r   r   test_atomic_cas3_2dim;  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_1dimA  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_2dimG  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_oldR  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   r5   r6   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   r5   rV  r  r   r   r   r   c  s    zBTestCudaAtomics.test_atomic_max_returns_no_replace.<locals>.kernelr  r  r  r   r   r   "test_atomic_max_returns_no_replaceb  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   j  s    zCTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernelr)   r  r  r   r   r   #test_atomic_max_returns_old_replacei  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   q  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_arrayp  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   r5   rV  r   r  r  r   r   r   r   x  s    zCTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernelr  r  r  r   r   r   #test_atomic_max_returns_old_nan_valw  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   r5   rr  r  r   r   r   r     s    zFTestCudaAtomics.test_atomic_min_returns_old_no_replace.<locals>.kernelr  r  r  r   r   r   &test_atomic_min_returns_old_no_replace~  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   r5   rr  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 )NrT  r   r)   r   r   )r   r   r   r   r   r   r   r   atomic_nanmaxnanmaxr   r   	r   r   rX  rY  init_valrZ  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 Nr]  r^  r   r   rX  rY  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 Nr]  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   r3   r   r)   r   r   rn  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   rm  r3   r   r)   r   rn  r   )r   r   r   r   r   r  r   r   r   ro  r  r   r   rj  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 )NrT  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   r}  r  r   r   rj  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   r5   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   r5   r  r   r  r  r   r   r   r   0  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   r5   r  r  r   r   r   r   7  s    zITestCudaAtomics.test_atomic_nanmin_returns_old_no_replace.<locals>.kernelr  r  r  r   r   r   )test_atomic_nanmin_returns_old_no_replace6  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   E  s    zKTestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_array.<locals>.kernelr  r  r   r   r   +test_atomic_nanmin_returns_old_nan_in_arrayD  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   r5   r  r   r  r  r   r   r   r   L  s    zFTestCudaAtomics.test_atomic_nanmin_returns_old_nan_val.<locals>.kernelr  r  r  r   r   r   &test_atomic_nanmin_returns_old_nan_valK  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/  r3  r5  r8  r;  r<  r=  r>  r?  r@  rA  rB  rC  rD  rE  rF  rG  rH  rI  rJ  rK  rL  rO  rQ  rR  rS  r[  r`  rb  rc  rd  re  rf  rk  rp  rs  rt  ru  rv  rw  rx  ry  r|  r~  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  __classcell__r   r   r   r   r     s&  					



	
	


	
	

r   __main__)qnumpyr   textwrapr   Znumbar   r   r   r   r   Znumba.cuda.testingr   r	   r
   Z
numba.corer   r   r   r   r   r&   r'   r.   r0   r1   r8   r9   r?   r@   rA   rF   rH   rK   rM   rN   rR   rS   rV   rW   rX   rZ   r[   r]   r^   r_   ra   rb   rc   rd   re   rf   rg   rh   rj   rk   rl   rm   rp   rr   rs   rv   rx   r{   r|   r}   r~   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   rU  ri  ro  r  rq  r{  r}  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   













	

	%           
