
    xKgQ                     .   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        Z ej                   d      d	        Z ej                   d      d
        Z ej                   d      d        Z ej                   d      d        Z ej                   d      d        Z ej                   d      d        Z ej                   d      d        Zd Zd Zd Zd Zd Zd Zd Zd Z d Z!d Z"d Z#d Z$d Z%d Z&d Z'd Z(d  Z)d! Z*d" Z+d# Z,d$ Z-d% Z.d& Z/d' Z0d( Z1d) Z2d* Z3d+ Z4d, Z5d- Z6d. Z7d/ Z8d0 Z9d1 Z:d2 Z;d3 Z<d4 Z=d5 Z>d6 Z?d7 Z@d8 ZAd9 ZBd: ZCd; ZDd< ZEd= ZFd> ZGd? ZHd@ ZIdA ZJdB ZKdC ZLdD ZMdE ZNdF ZOdG ZPdH ZQdI ZRdJ ZSdK ZTdL ZUdM ZVdN ZWdO ZXdP ZYdQ ZZ eZdR      \  Z[Z\Z]Z^ eZdS      \  Z_Z`ZaZb eZdT      \  ZcZdZeZf eZdU      \  ZgZhZiZjdV ZkdW ZldX Zm G dY dZe      Zneod[k(  r ej                          yy)\    N)dedent)cudauint32uint64float32float64)unittestCUDATestCasecc_X_or_above)configT)devicec                     t        |       S N)r   nums    h/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_atomics.pyatomic_cast_to_uint64r   	   s    #;    c                     t        |       S r   )intr   s    r   atomic_cast_to_intr      s    s8Or   c                     | S r    r   s    r   atomic_cast_noner      s    Jr   c	                     t         j                  j                  }	t         j                  j	                  ||      }
||
|	<   t        j
                           |||	   |z        }|r||z  } ||
||       t        j
                          |
|	   | |	<   y r   r   	threadIdxxsharedarraysyncthreads)aryidxop2	ary_dtypeary_nelements
binop_func	cast_funcinitializerneg_idxtidsmbins               r   atomic_binary_1dim_sharedr.      s     ..

C			=)	4BBsG
CH},
-CM!r3#wCHr   c                    t         j                  j                  }t         j                  j	                  ||      }| |   ||<   t        j
                           |||   |z        }	 |||	|       t        j
                          ||   | |<   y r   r   )
r"   r#   r$   r%   r&   r'   r(   r+   r,   r-   s
             r   atomic_binary_1dim_shared2r0   (   sz     ..

C			=)	4B#hBsG
CH},
-Cr3#wCHr   c                    t         j                  j                  }t         j                  j                  }t         j                  j                  ||      }	| ||f   |	||f<   t        j                          | ||      f}
|r|
d   |d   z  |
d   |d   z  f}
 ||	|
|       t        j                          |	||f   | ||f<   y Nr      )r   r   r   yr   r    r!   )r"   r$   r%   	ary_shaper'   y_cast_funcr*   txtyr,   r-   s              r   atomic_binary_2dim_sharedr9   5   s     
		B			B			9i	0BRVBr2vJ{2
C1v	!$c!fy|&;<r3RV*CBKr   c                     t         j                  j                  }t         j                  j                  }| ||      f}|r*|d   | j                  d   z  |d   | j                  d   z  f} || ||       y r2   )r   r   r   r4   shape)r"   r$   r'   r6   r*   r7   r8   r-   s           r   atomic_binary_2dim_globalr<   E   si    			B			B{2
C1v		!$c!fsyy|&;<sCr   c                 |    t         j                  j                  }t        ||   |z        }|r||z  } || ||       y r   )r   r   r   r   )r"   r#   r&   r$   r'   r*   r+   r-   s           r   atomic_binary_1dim_globalr>   O   s?     ..

C
c#h&
'CM!sCr   c                 j    t        | | dt        dt        j                  j                  t
        dd	       y Nr3       r   Fr.   r   r   atomicaddr   r"   s    r   
atomic_addrF   Y   '    c362"kkoo/?EKr   c                 j    t        | | dt        dt        j                  j                  t
        dd	       y )Nr3   rA   r   TrB   rE   s    r   atomic_add_wraprI   ^   s'    c362"kkoo/?DJr   c           	      f    t        | dt        dt        j                  j                  t
        d       y Nr3         Fr9   r   r   rC   rD   r   rE   s    r   atomic_add2rP   c   #    c1ff"kkoo/?Hr   c           	      f    t        | dt        dt        j                  j                  t
        d       y )Nr3   rL   TrO   rE   s    r   atomic_add2_wraprS   h   s#    c1ff"kkoo/?Gr   c           	      f    t        | dt        dt        j                  j                  t
        d       y rK   )r9   r   r   rC   rD   r   rE   s    r   atomic_add3rU   m   #    c1ff"kkoo/DeMr   c                 j    t        | | dt        dt        j                  j                  t
        dd	       y N      ?rA           Fr.   r   r   rC   rD   r   rE   s    r   atomic_add_floatr\   r   '    c3Wb"kkoo/A3Or   c                 j    t        | | dt        dt        j                  j                  t
        dd	       y NrY   rA   rZ   Tr[   rE   s    r   atomic_add_float_wrapr`   w   s'    c3Wb"kkoo/A3Nr   c           	      f    t        | dt        dt        j                  j                  t
        d       y NrY   rL   Fr9   r   r   rC   rD   r   rE   s    r   atomic_add_float_2rd   |   #    c3"kkoo/?Hr   c           	      f    t        | dt        dt        j                  j                  t
        d       y NrY   rL   Trc   rE   s    r   atomic_add_float_2_wraprh      #    c3"kkoo/?Gr   c           	      f    t        | dt        dt        j                  j                  t
        d       y rb   )r9   r   r   rC   rD   r   rE   s    r   atomic_add_float_3rk      #    c3"kkoo/DeMr   c                 T    t        || ddt        j                  j                  d       y NrA   rY   Fr>   r   rC   rD   r#   r"   s     r   atomic_add_double_globalrq          c3C%Hr   c                 T    t        || ddt        j                  j                  d       y )NrA   rY   Tro   rp   s     r   atomic_add_double_global_wraprt      s    c3C$Gr   c                 Z    t        | dt        j                  j                  t        d       y Nr3   Fr<   r   rC   rD   r   rE   s    r   atomic_add_double_global_2rx      s    c1dkkoo7GOr   c                 Z    t        | dt        j                  j                  t        d       y )Nr3   Trw   rE   s    r   atomic_add_double_global_2_wraprz      s    c1dkkoo7GNr   c                 Z    t        | dt        j                  j                  t        d       y rv   )r<   r   rC   rD   r   rE   s    r   atomic_add_double_global_3r|      s    c1dkkoo7L#%r   c                 j    t        || dt        dt        j                  j                  t
        dd	       y rX   r.   r   r   rC   rD   r   rp   s     r   atomic_add_doubler      '    c3Wb"kkoo/?eMr   c                 j    t        || dt        dt        j                  j                  t
        dd	       y r_   r~   rp   s     r   atomic_add_double_wrapr      s'    c3Wb"kkoo/?dLr   c           	      f    t        | dt        dt        j                  j                  t
        d       y rb   r9   r   r   rC   rD   r   rE   s    r   atomic_add_double_2r      re   r   c           	      f    t        | dt        dt        j                  j                  t
        d       y rg   r   rE   s    r   atomic_add_double_2_wrapr      ri   r   c           	      f    t        | dt        dt        j                  j                  t
        d       y rb   )r9   r   r   rC   rD   r   rE   s    r   atomic_add_double_3r      rl   r   c                 j    t        | | dt        dt        j                  j                  t
        dd	       y r@   )r.   r   r   rC   subr   rE   s    r   
atomic_subr      rG   r   c           	      f    t        | dt        dt        j                  j                  t
        d       y rK   )r9   r   r   rC   r   r   rE   s    r   atomic_sub2r      rQ   r   c           	      f    t        | dt        dt        j                  j                  t
        d       y rK   )r9   r   r   rC   r   r   rE   s    r   atomic_sub3r      rV   r   c                 j    t        | | dt        dt        j                  j                  t
        dd	       y rX   )r.   r   r   rC   r   r   rE   s    r   atomic_sub_floatr      r]   r   c           	      f    t        | dt        dt        j                  j                  t
        d       y rb   )r9   r   r   rC   r   r   rE   s    r   atomic_sub_float_2r      re   r   c           	      f    t        | dt        dt        j                  j                  t
        d       y rb   )r9   r   r   rC   r   r   rE   s    r   atomic_sub_float_3r      rl   r   c                 j    t        || dt        dt        j                  j                  t
        dd	       y rX   )r.   r   r   rC   r   r   rp   s     r   atomic_sub_doubler      r   r   c           	      f    t        | dt        dt        j                  j                  t
        d       y rb   )r9   r   r   rC   r   r   rE   s    r   atomic_sub_double_2r      re   r   c           	      f    t        | dt        dt        j                  j                  t
        d       y rb   r9   r   r   rC   r   r   rE   s    r   atomic_sub_double_3r      rl   r   c                 T    t        || ddt        j                  j                  d       y rn   )r>   r   rC   r   rp   s     r   atomic_sub_double_globalr      rr   r   c                 Z    t        | dt        j                  j                  t        d       y )NrY   F)r<   r   rC   r   r   rE   s    r   atomic_sub_double_global_2r      s    c39I#%r   c           	      f    t        | dt        dt        j                  j                  t
        d       y rb   r   rE   s    r   atomic_sub_double_global_3r      rl   r   c                 j    t        | | |t        dt        j                  j                  t
        dd	       y )NrA   r3   F)r.   r   r   rC   and_r   r"   r$   s     r   
atomic_andr      s)    c3VR"kk..0@!ULr   c           	      f    t        | |t        dt        j                  j                  t
        d       y NrL   F)r9   r   r   rC   r   r   r   s     r   atomic_and2r      %    c3"kk..0@%Ir   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   atomic_and3r      s%    c3"kk..0EuNr   c                 T    t        || d|t        j                  j                  d       y NrA   F)r>   r   rC   r   r#   r"   r$   s      r   atomic_and_globalr         c3C1A1A5Ir   c                 Z    t        | |t        j                  j                  t        d       y NF)r<   r   rC   r   r   r   s     r   atomic_and_global_2r     s    c3(8(8.7r   c                 j    t        | | |t        dt        j                  j                  t
        dd	       y NrA   r   F)r.   r   r   rC   or_r   r   s     r   	atomic_orr     '    c3VR"kkoo/?EKr   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   
atomic_or2r     #    c3"kkoo/?Hr   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   
atomic_or3r     #    c3"kkoo/DeMr   c                 T    t        || d|t        j                  j                  d       y r   )r>   r   rC   r   r   s      r   atomic_or_globalr     rr   r   c                 Z    t        | |t        j                  j                  t        d       y r   )r<   r   rC   r   r   r   s     r   atomic_or_global_2r         c3.7r   c                 j    t        | | |t        dt        j                  j                  t
        dd	       y r   )r.   r   r   rC   xorr   r   s     r   
atomic_xorr   $  r   r   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   atomic_xor2r   )  r   r   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   atomic_xor3r   .  r   r   c                 T    t        || d|t        j                  j                  d       y r   )r>   r   rC   r   r   s      r   atomic_xor_globalr   3  rr   r   c                 Z    t        | |t        j                  j                  t        d       y r   )r<   r   rC   r   r   r   s     r   atomic_xor_global_2r   7  r   r   c           	      f    t        | ||t        dt        j                  j                  t
               y NrA   )r0   r   r   rC   incr   r"   r#   r$   s      r   atomic_inc32r   <  #    sCfb#{{0@Br   c           	      f    t        | ||t        dt        j                  j                  t
               y r   )r0   r   r   rC   r   r   r   s      r   atomic_inc64r   A  #    sCfb#{{0BDr   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   atomic_inc2_32r   F  r   r   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   atomic_inc2_64r   K  r   r   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   atomic_inc3r   P  r   r   c                 T    t        || d|t        j                  j                  d       y r   )r>   r   rC   r   r   s      r   atomic_inc_globalr   U  rr   r   c                 Z    t        | |t        j                  j                  t        d       y r   )r<   r   rC   r   r   r   s     r   atomic_inc_global_2r   Y  r   r   c           	      f    t        | ||t        dt        j                  j                  t
               y r   )r0   r   r   rC   decr   r   s      r   atomic_dec32r   ^  r   r   c           	      f    t        | ||t        dt        j                  j                  t
               y r   )r0   r   r   rC   r   r   r   s      r   atomic_dec64r   c  r   r   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   atomic_dec2_32r   h  r   r   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   atomic_dec2_64r   m  r   r   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   atomic_dec3r   r  r   r   c                 T    t        || d|t        j                  j                  d       y r   )r>   r   rC   r   r   s      r   atomic_dec_globalr   w  rr   r   c                 Z    t        | |t        j                  j                  t        d       y r   )r<   r   rC   r   r   r   s     r   atomic_dec_global_2r   {  r   r   c           	      f    t        | ||t        dt        j                  j                  t
               y r   )r0   r   r   rC   exchr   r   s      r   atomic_exchr     s%    sCfb#{{//1ACr   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   atomic_exch2r     r   r   c           	      f    t        | |t        dt        j                  j                  t
        d       y r   )r9   r   r   rC   r   r   r   s     r   atomic_exch3r     r   r   c                 T    t        || d|t        j                  j                  d       y r   )r>   r   rC   r   r   s      r   atomic_exch_globalr     r   r   c                     t        d      j                  |       }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   rC   atomic_double_normalizedindexatomic_double_oneindexatomic_double_shared)r   formatexecr   r   r   )r   fnslds      r   gen_atomic_extreme_funcsr    sh    
  	6 
T	7 8 
Bt6BBGxL"<='("-C*DF Fr   zcuda.atomic.maxzcuda.atomic.minzcuda.atomic.nanmaxzcuda.atomic.nanminc                     t        j                  d      }|| j                  k  r+t         j                  j	                  | |d  |||         ||<   y y Nr3   )r   gridsizerC   compare_and_swapresoldr"   fill_valgids        r   atomic_compare_and_swapr    sE    
))A,C
SXX~;;//CD	8SXNC r   c                     t        j                  d      }|| j                  k  r)t         j                  j	                  | ||||         ||<   y y r  )r   r  r	  rC   casr  s        r   atomic_cas_1dimr    s?    
))A,C
SXX~;;??3Xs3x@C r   c                     t        j                  d      }|d   | j                  d   k  r?|d   | j                  d   k  r)t         j                  j	                  | ||||         ||<   y y y )N   r   r3   )r   r  r;   rC   r  r  s        r   atomic_cas_2dimr    s^    
))A,C
1v		!Q#))A,!6;;??3Xs3x@C "7r   c                       e Zd Z fdZd Zd Zd Zd Zd Zd Z	ddZ
d	 Zd
 Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Z d Z!d  Z"d! Z#d" Z$d# Z%d$ Z&d% Z'd& Z(d' Z)d( Z*d) Z+d* Z,d+ Z-d, Z.d- Z/d. Z0d/ Z1d0 Z2d1 Z3d2 Z4d3 Z5d4 Z6d5 Z7d6 Z8d7 Z9d8 Z:d9 Z;d: Z<d; Z=d< Z>d= Z?d> Z@d? ZAd@ ZBdA ZCdB ZDdC ZEdD ZFdE ZGdF ZHdG ZIdH ZJdI ZKdJ ZLdK ZMdL ZNdM ZOdN ZPdO ZQdP ZRdQ ZSdR ZTdS ZUdT ZVdU ZWdV ZXdW ZYdX ZZdY Z[dZ Z\d[ Z]d\ Z^d] Z_d^ Z`d_ Zad` Zbda ZcddbZddc Zedd Zfde Zgdf Zhdg Zidh Zjdi Zkdj Zldk Zmdl Zndm Zodn Zpdo Zqdp Zrdq Zsdr Ztds Zudt Zvdu Zwdv Zxdw Zydx Zzdy Z{dz Z|d{ Z}d| Z~d} Zd~ Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Z xZS )TestCudaAtomicsc                 `    t         |           t        j                  j	                  d       y )Nr   )supersetUpnprandomseed)self	__class__s    r   r  zTestCudaAtomics.setUp  s    
		qr   c                    t         j                  j                  ddd      j                  t         j                        }|j                         }|j                         } t        j                  d      t              } |d   |        t        j                  d      t              } |d   |       t        j                  dt         j                        }t        |j                        D ]  }|||   xx   dz  cc<    | j                  t        j                  ||k(               | j                  t        j                  ||k(               y Nr   rA   r	  zvoid(uint32[:])r3   rA   dtyper3   )r  r  randintastyper   copyr   jitrF   rI   zerosranger	  
assertTrueall)r  r"   ary_wraporigcuda_atomic_addcuda_atomic_add_wrapgoldis           r   test_atomic_addzTestCudaAtomics.test_atomic_add  s    ii2B/66ryyA88:xxz5$((#45jAs#:txx(9:?K#U#H-xx")),tyy!AaMQM " 	sd{+,x4/01r   c                 2   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         }|j                         } t        j                  d      t              } |d   |        t        j                  d      t              } |d   |       | j                  t        j                  ||dz   k(               | j                  t        j                  ||dz   k(               y 	Nr   rA   r#  rM   rN   zvoid(uint32[:,:])r3   rL   r3   )r  r  r'  r(  r   reshaper)  r   r*  rP   rS   r-  r.  )r  r"   r/  r0  cuda_atomic_add2cuda_atomic_add2_wraps         r   test_atomic_add2z TestCudaAtomics.test_atomic_add2  s    ii2B/66ryyAII!QO88:xxz8488$78E##C( =)< =>N O(i(2sdQh/0x4!8345r   c                 j   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       | j                  t        j                  ||dz   k(               y r7  )r  r  r'  r(  r   r9  r)  r   r*  rU   r-  r.  r  r"   r0  cuda_atomic_add3s       r   test_atomic_add3z TestCudaAtomics.test_atomic_add3  s    ii2B/66ryyAII!QOxxz8488$78E##C(sdQh/0r   c                    t         j                  j                  ddd      j                  t         j                        }|j                         }|j                         j                  t         j                        } t        j                  d      t              } |d   |        t        j                  d      t              } |d   |       t        j                  dt         j                        }t        |j                        D ]  }|||   xx   dz  cc<    | j                  t        j                   ||k(               | j                  t        j                   ||k(               y Nr   rA   r#  zvoid(float32[:])r$  r%  rY   )r  r  r'  r(  r   r)  intpr   r*  r\   r`   r+  r   r,  r	  r-  r.  )r  r"   r/  r0  cuda_atomic_add_floatadd_float_wrapr3  r4  s           r   test_atomic_add_floatz%TestCudaAtomics.test_atomic_add_float  s   ii2B/66rzzB88:xxz  ) <); <=M N$e$S)5"456KLuh'xx")),tyy!AaMS M " 	sd{+,x4/01r   c                 2   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         }|j                         } t        j                  d      t              } |d   |        t        j                  d      t              } |d   |       | j                  t        j                  ||dz   k(               | j                  t        j                  ||dz   k(               y 	Nr   rA   r#  rM   rN   zvoid(float32[:,:])r8  r3   )r  r  r'  r(  r   r9  r)  r   r*  rd   rh   r-  r.  )r  r"   r/  r0  r:  cuda_func_wraps         r   test_atomic_add_float_2z'TestCudaAtomics.test_atomic_add_float_2  s    ii2B/66rzzBJJ1aP88:xxz9488$89:LM##C(7"678OP!y!(+sdQh/0x4!8345r   c                 j   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       | j                  t        j                  ||dz   k(               y rH  )r  r  r'  r(  r   r9  r)  r   r*  rk   r-  r.  r>  s       r   test_atomic_add_float_3z'TestCudaAtomics.test_atomic_add_float_3"  s    ii2B/66rzzBJJ1aPxxz9488$89:LM##C(sdQh/0r   c                 r   t         j                  ry t        t        |j	                         j                                     }t        dd      rCt        j                  j                         dkD  rd}nd}|r| d}| j                  | d|       y |r| j                  d|       y | j                  d	|       y )
N   r   )   r3   redatomz.sharedz.add.f64zatom.shared.cas.b64zatom.cas.b64)r   ENABLE_CUDASIMnextiterinspect_asmvaluesr   r   runtimeget_versionassertIn)r  kernelr   asminsts        r   assertCorrectFloat64Atomicsz+TestCudaAtomics.assertCorrectFloat64Atomics*  s       4**,33567A||'')G3 w'MMTF(+S13S9nc2r   c                    t         j                  j                  dddt         j                        }t        j                  dt         j
                        }|j                         } t        j                  d      t              } |d   ||        t        j                  d      t              } |d   ||       t        j                  dt         j                        }t        |j                        D ]  }|||   xx   dz  cc<    t         j                  j                  ||       t         j                  j                  ||       | j!                  |       | j!                  |       y Nr   rA   r	  r&  void(int64[:], float64[:])r$  r%  rY   )r  r  r'  int64r+  r   r)  r   r*  r   r   r   r,  r	  testingassert_equalr]  )r  r#   r"   r/  cuda_fnwrap_fnr3  r4  s           r   test_atomic_add_doublez&TestCudaAtomics.test_atomic_add_doubleC  s   ii2Bbhh?hhr2::&88:8$((789JKsC 8$((789OPsH%xx")),sxxAQLCL ! 	

T*


$/((1((1r   c                 Z   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         }|j                         } t        j                  d      t              } |d   |        t        j                  d      t              } |d   |       t         j                  j                  ||dz          t         j                  j                  ||dz          | j                  |       | j                  |       y 	Nr   rA   r#  rM   rN   void(float64[:,:])r8  r3   )r  r  r'  r(  r   r9  r)  r   r*  r   r   rc  rd  r]  )r  r"   r/  r0  re  cuda_fn_wraps         r   test_atomic_add_double_2z(TestCudaAtomics.test_atomic_add_double_2W  s    ii2B/66rzzBJJ1aP88:xxz0$((/01DE	35txx 456NOY)


TAX.


$(3((1((6r   c                 ~   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       t         j                  j                  ||dz          | j                  |       y ri  )r  r  r'  r(  r   r9  r)  r   r*  r   rc  rd  r]  r  r"   r0  	cuda_funcs       r   test_atomic_add_double_3z(TestCudaAtomics.test_atomic_add_double_3g  s    ii2B/66rzzBJJ1aPxxz2DHH123FG		)S!


TAX.((3r   c                    t         j                  j                  dddt         j                        }t        j                  dt         j
                        }|j                         }d} t        j                  |      t              } t        j                  |      t              } |d   ||        |d   ||       t        j                  dt         j                        }t        |j                        D ]  }|||   xx   dz  cc<    t         j                  j                  ||       t         j                  j                  ||       | j!                  |d	       | j!                  |d	       y )
Nr   rA   r`  ra  r$  r%  rY   Fr   )r  r  r'  rb  r+  r   r)  r   r*  rq   rt   r   r,  r	  rc  rd  r]  )	r  r#   r"   r/  sigro  wrap_cuda_funcr3  r4  s	            r   test_atomic_add_double_globalz-TestCudaAtomics.test_atomic_add_double_globalp  s   ii2Bbhh?hhr2::&88:*!DHHSM":;	&#'DE	%c"uc8,xx")),sxxAQLCL ! 	

T*


$/((5(A(((Fr   c                 f   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         }|j                         }d} t        j                  |      t              } t        j                  |      t              } |d   |        |d   |       t         j                  j                  ||dz          t         j                  j                  ||dz          | j                  |d	
       | j                  |d	
       y Nr   rA   r#  rM   rN   rj  r8  r3   Frr  )r  r  r'  r(  r   r9  r)  r   r*  rx   rz   rc  rd  r]  )r  r"   r/  r0  rs  ro  rt  s          r   test_atomic_add_double_global_2z/TestCudaAtomics.test_atomic_add_double_global_2  s    ii2B/66rzzBJJ1aP88:xxz"!DHHSM"<=	&#'FG	)S!!y!(+


TAX.


$(3((5(A(((Fr   c                    t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       t         j                  j                  ||dz          | j                  |d	
       y rw  )r  r  r'  r(  r   r9  r)  r   r*  r|   rc  rd  r]  rn  s       r   test_atomic_add_double_global_3z/TestCudaAtomics.test_atomic_add_double_global_3  s    ii2B/66rzzBJJ1aPxxz2DHH123MN		)S!


TAX.((5(Ar   c                    t         j                  j                  ddd      j                  t         j                        }|j                         } t        j                  d      t              } |d   |       t        j                  dt         j                        }t        |j                        D ]  }|||   xx   dz  cc<    | j                  t        j                  ||k(               y r"  )r  r  r'  r(  r   r)  r   r*  r   r+  r,  r	  r-  r.  )r  r"   r0  cuda_atomic_subr3  r4  s         r   test_atomic_subzTestCudaAtomics.test_atomic_sub  s    ii2B/66ryyAxxz5$((#45jAs#xx")),tyy!AaMQM " 	sd{+,r   c                 j   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       | j                  t        j                  ||dz
  k(               y r7  )r  r  r'  r(  r   r9  r)  r   r*  r   r-  r.  r  r"   r0  cuda_atomic_sub2s       r   test_atomic_sub2z TestCudaAtomics.test_atomic_sub2      ii2B/66ryyAII!QOxxz8488$78E##C(sdQh/0r   c                 j   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       | j                  t        j                  ||dz
  k(               y r7  )r  r  r'  r(  r   r9  r)  r   r*  r   r-  r.  r  r"   r0  cuda_atomic_sub3s       r   test_atomic_sub3z TestCudaAtomics.test_atomic_sub3  r  r   c                    t         j                  j                  ddd      j                  t         j                        }|j                         j                  t         j                        } t        j                  d      t              } |d   |       t        j                  dt         j                        }t        |j                        D ]  }|||   xx   dz  cc<    | j                  t        j                  ||k(               y rB  )r  r  r'  r(  r   r)  rC  r   r*  r   r+  r,  r	  r-  r.  )r  r"   r0  cuda_atomic_sub_floatr3  r4  s         r   test_atomic_sub_floatz%TestCudaAtomics.test_atomic_sub_float  s    ii2B/66rzzBxxz  ) <); <=M N$e$S)xx"**-tyy!AaMS M " 	sd{+,r   c                 j   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       | j                  t        j                  ||dz
  k(               y rH  )r  r  r'  r(  r   r9  r)  r   r*  r   r-  r.  r  s       r   test_atomic_sub_float_2z'TestCudaAtomics.test_atomic_sub_float_2      ii2B/66rzzBJJ1aPxxz9488$89:LM##C(sdQh/0r   c                 j   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       | j                  t        j                  ||dz
  k(               y rH  )r  r  r'  r(  r   r9  r)  r   r*  r   r-  r.  r  s       r   test_atomic_sub_float_3z'TestCudaAtomics.test_atomic_sub_float_3  r  r   c                    t         j                  j                  dddt         j                        }t        j                  dt         j
                        } t        j                  d      t              } |d   ||       t        j                  dt         j
                        }t        |j                        D ]  }|||   xx   dz  cc<    t         j                  j                  ||       y r_  )r  r  r'  rb  r+  r   r   r*  r   r,  r	  rc  rd  )r  r#   r"   ro  r3  r4  s         r   test_atomic_sub_doublez&TestCudaAtomics.test_atomic_sub_double  s    ii2Bbhh?hhr2::&:DHH9:;LM		%c"xx"**-sxxAQLCL ! 	

T*r   c                 \   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       t         j                  j                  ||dz
         y ri  )r  r  r'  r(  r   r9  r)  r   r*  r   rc  rd  rn  s       r   test_atomic_sub_double_2z(TestCudaAtomics.test_atomic_sub_double_2      ii2B/66rzzBJJ1aPxxz2DHH123FG		)S!


TAX.r   c                 \   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       t         j                  j                  ||dz
         y ri  )r  r  r'  r(  r   r9  r)  r   r*  r   rc  rd  rn  s       r   test_atomic_sub_double_3z(TestCudaAtomics.test_atomic_sub_double_3  r  r   c                    t         j                  j                  dddt         j                        }t        j                  dt         j
                        }d} t        j                  |      t              } |d   ||       t        j                  dt         j
                        }t        |j                        D ]  }|||   xx   dz  cc<    t         j                  j                  ||       y r_  )r  r  r'  rb  r+  r   r   r*  r   r,  r	  rc  rd  )r  r#   r"   rs  ro  r3  r4  s          r   test_atomic_sub_double_globalz-TestCudaAtomics.test_atomic_sub_double_global  s    ii2Bbhh?hhr2::&*!DHHSM":;		%c"xx"**-sxxAQLCL ! 	

T*r   c                 \   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       t         j                  j                  ||dz
         y ri  )r  r  r'  r(  r   r9  r)  r   r*  r   rc  rd  rn  s       r   test_atomic_sub_double_global_2z/TestCudaAtomics.test_atomic_sub_double_global_2      ii2B/66rzzBJJ1aPxxz2DHH123MN		)S!


TAX.r   c                 \   t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   |       t         j                  j                  ||dz
         y ri  )r  r  r'  r(  r   r9  r)  r   r*  r   rc  rd  rn  s       r   test_atomic_sub_double_global_3z/TestCudaAtomics.test_atomic_sub_double_global_3  r  r   c                    t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        }|j                         } t        j                  d      t              } |d   ||       |j                         }t        |j                        D ]  }|||   xx   |z  cc<    | j                  t        j                  ||k(               y )N  r   rA   r#  void(uint32[:], uint32)r$  )r  r  r'  r(  r   r)  r   r*  r   r,  r	  r-  r.  r  
rand_constr"   r0  ro  r3  r4  s          r   test_atomic_andzTestCudaAtomics.test_atomic_and  s    YY&&s+
ii2B/66ryyAxxz7DHH67
C		%j)xxztyy!AaMZ'M " 	sd{+,r   c                    t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   ||       | j                  t        j                  |||z  k(               y 	Nr  r   rA   r#  rM   rN   void(uint32[:,:], uint32)r8  )r  r  r'  r(  r   r9  r)  r   r*  r   r-  r.  r  r  r"   r0  cuda_atomic_and2s        r   test_atomic_and2z TestCudaAtomics.test_atomic_and2      YY&&s+
ii2B/66ryyAII!QOxxz@488$?@M##C4sdZ&7789r   c                    t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   ||       | j                  t        j                  |||z  k(               y r  )r  r  r'  r(  r   r9  r)  r   r*  r   r-  r.  r  r  r"   r0  cuda_atomic_and3s        r   test_atomic_and3z TestCudaAtomics.test_atomic_and3  r  r   c                    t         j                  j                  d      }t         j                  j                  dddt         j                        }t         j                  j                  dddt         j                        }d} t	        j
                  |      t              } |d   |||       |j                         }t        |j                        D ]  }|||   xx   |z  cc<    t         j                  j                  ||       y Nr  r   rA   r`  zvoid(int32[:], int32[:], int32)r$  )r  r  r'  int32r   r*  r   r)  r,  r	  rc  rd  r  r  r#   r"   rs  ro  r3  r4  s           r   test_atomic_and_globalz&TestCudaAtomics.test_atomic_and_global%  s    YY&&s+
ii2Bbhh?ii2Bbhh?/!DHHSM"34		%c:.xxzsxxAQLJ&L ! 	

T*r   c                    t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   ||       t         j                  j                  |||z         y r  )r  r  r'  r(  r   r9  r)  r   r*  r   rc  rd  r  r  r"   r0  ro  s        r   test_atomic_and_global_2z(TestCudaAtomics.test_atomic_and_global_23      YY&&s+
ii2B/66ryyAII!QOxxz9DHH89:MN		)S*-


TJ%67r   c                 "   t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        }|j                         } t        j                  d      t              } |d   ||       t        j                  dt         j                        }t        |j                        D ]  }|||   xx   |z  cc<    | j                  t        j                  ||k(               y Nr  r   rA   r#  r  r$  r%  )r  r  r'  r(  r   r)  r   r*  r   r+  r,  r	  r-  r.  r  s          r   test_atomic_orzTestCudaAtomics.test_atomic_or;  s    YY&&s+
ii2B/66ryyAxxz7DHH67	B		%j)xx")),tyy!AaMZ'M " 	sd{+,r   c                    t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   ||       | j                  t        j                  |||z  k(               y r  )r  r  r'  r(  r   r9  r)  r   r*  r   r-  r.  r  s        r   test_atomic_or2zTestCudaAtomics.test_atomic_or2H      YY&&s+
ii2B/66ryyAII!QOxxz@488$?@L##C4sdZ&7789r   c                    t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   ||       | j                  t        j                  |||z  k(               y r  )r  r  r'  r(  r   r9  r)  r   r*  r   r-  r.  r  s        r   test_atomic_or3zTestCudaAtomics.test_atomic_or3P  r  r   c                    t         j                  j                  d      }t         j                  j                  dddt         j                        }t         j                  j                  dddt         j                        }d} t	        j
                  |      t              } |d   |||       |j                         }t        |j                        D ]  }|||   xx   |z  cc<    t         j                  j                  ||       y r  )r  r  r'  r  r   r*  r   r)  r,  r	  rc  rd  r  s           r   test_atomic_or_globalz%TestCudaAtomics.test_atomic_or_globalX  s    YY&&s+
ii2Bbhh?ii2Bbhh?/!DHHSM"23		%c:.xxzsxxAQLJ&L ! 	

T*r   c                    t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   ||       t         j                  j                  |||z         y r  )r  r  r'  r(  r   r9  r)  r   r*  r   rc  rd  r  s        r   test_atomic_or_global_2z'TestCudaAtomics.test_atomic_or_global_2f  s    YY&&s+
ii2B/66ryyAII!QOxxz9DHH89:LM		)S*-


TJ%67r   c                 "   t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        }|j                         } t        j                  d      t              } |d   ||       t        j                  dt         j                        }t        |j                        D ]  }|||   xx   |z  cc<    | j                  t        j                  ||k(               y r  )r  r  r'  r(  r   r)  r   r*  r   r+  r,  r	  r-  r.  r  s          r   test_atomic_xorzTestCudaAtomics.test_atomic_xorn  s    YY&&s+
ii2B/66ryyAxxz7DHH67
C		%j)xx")),tyy!AaMZ'M " 	sd{+,r   c                    t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   ||       | j                  t        j                  |||z  k(               y r  )r  r  r'  r(  r   r9  r)  r   r*  r   r-  r.  )r  r  r"   r0  cuda_atomic_xor2s        r   test_atomic_xor2z TestCudaAtomics.test_atomic_xor2{  r  r   c                    t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   ||       | j                  t        j                  |||z  k(               y r  )r  r  r'  r(  r   r9  r)  r   r*  r   r-  r.  )r  r  r"   r0  cuda_atomic_xor3s        r   test_atomic_xor3z TestCudaAtomics.test_atomic_xor3  r  r   c                    t         j                  j                  d      }t         j                  j                  dddt         j                        }t         j                  j                  dddt         j                        }|j	                         }d} t        j                  |      t              } |d   |||       t        |j                        D ]  }|||   xx   |z  cc<    t         j                  j                  ||       y r  )r  r  r'  r  r)  r   r*  r   r,  r	  rc  rd  )r  r  r#   r"   r3  rs  ro  r4  s           r   test_atomic_xor_globalz&TestCudaAtomics.test_atomic_xor_global  s    YY&&s+
ii2Bbhh?ii2Bbhh?xxz/!DHHSM"34		%c:.sxxAQLJ&L ! 	

T*r   c                    t         j                  j                  d      }t         j                  j                  ddd      j                  t         j                        j                  dd      }|j                         } t        j                  d      t              } |d   ||       t         j                  j                  |||z         y r  )r  r  r'  r(  r   r9  r)  r   r*  r   rc  rd  r  s        r   test_atomic_xor_global_2z(TestCudaAtomics.test_atomic_xor_global_2  r  r   c                     t         j                  j                  d|      }t         j                  j                  ddd      j                  |      }t        j                  d|      }|||fS )NrA   r%  r   r#  )r  r  r'  r(  arange)r  r&  rconstraryary_idxs        r   inc_dec_1dim_setupz"TestCudaAtomics.inc_dec_1dim_setup  s^    ""2e"4yy  BR 077>))Be,tW$$r   c                     t         j                  j                  d|      }t         j                  j                  ddd      j                  |      j	                  dd      }||fS )NrA   r%  r   r#  rM   rN   )r  r  r'  r(  r9  )r  r&  r  r  s       r   inc_dec_2dim_setupz"TestCudaAtomics.inc_dec_2dim_setup  sX    ""2U"3yy  BR 077>FFq!Lt|r   c           	          |j                         } t        j                  |      |      }	 |	||f   |||       t        j                  j                  |t        j                  ||k\  d|dz                y r2   r)  r   r*  r  rc  rd  where
r  r"   r#   r  rs  nblocksblksizer   r0  ro  s
             r   check_inc_indexzTestCudaAtomics.check_inc_index  b    xxz!DHHSM$'	#	'7"#Cf5


RXXdfna%JKr   c           	          |j                         } t        j                  |      |      }	 |	||f   |||       t        j                  j                  |t        j                  ||k\  d|dz                y r2   r  r  s
             r   check_inc_index2z TestCudaAtomics.check_inc_index2  r  r   c           	          |j                         } t        j                  |      |      } |||f   ||       t        j                  j                  |t        j                  ||k\  d|dz                y r2   r  	r  r"   r  rs  r  r  r   r0  ro  s	            r   	check_inczTestCudaAtomics.check_inc  s`    xxz!DHHSM$'	#	'7"#C0


RXXdfna%JKr   c           	          | j                  t        j                        \  }}}d}| j                  ||||ddt               y Nr%  "void(uint32[:], uint32[:], uint32)r3   rA   )r  r  r   r  r   r  r  r"   r#   rs  s        r   test_atomic_inc_32z"TestCudaAtomics.test_atomic_inc_32  ?    #66RYY6G
C2S#z32|Lr   c           	          | j                  t        j                        \  }}}d}| j                  ||||ddt               y Nr%  z"void(uint64[:], uint64[:], uint64)r3   rA   )r  r  r   r  r   r  s        r   test_atomic_inc_64z"TestCudaAtomics.test_atomic_inc_64  r  r   c                     | j                  t        j                        \  }}d}| j                  |||ddt               y Nr  r3   rL   )r  r  r   r  r   r  r  r"   rs  s       r   test_atomic_inc2_32z#TestCudaAtomics.test_atomic_inc2_32  5    11"))<
C)sJQ~Fr   c                     | j                  t        j                        \  }}d}| j                  |||ddt               y Nvoid(uint64[:,:], uint64)r3   rL   )r  r  r   r  r   r  s       r   test_atomic_inc2_64z#TestCudaAtomics.test_atomic_inc2_64  r  r   c                     | j                  t        j                        \  }}d}| j                  |||ddt               y r  )r  r  r   r  r   r  s       r   test_atomic_inc3z TestCudaAtomics.test_atomic_inc3  5    11"))<
C)sJQ{Cr   c           	          | j                  t        j                        \  }}}d}| j                  ||||ddt               y r  )r  r  r   r  r   r  s        r   test_atomic_inc_global_32z)TestCudaAtomics.test_atomic_inc_global_32  B    #66RYY6G
C2c3
CB/	1r   c           	          | j                  t        j                        \  }}}d}| j                  ||||ddt               y r  )r  r  r   r  r   r  s        r   test_atomic_inc_global_64z)TestCudaAtomics.test_atomic_inc_global_64  r  r   c                     | j                  t        j                        \  }}d}| j                  |||ddt               y r  )r  r  r   r  r   r  s       r   test_atomic_inc_global_2_32z+TestCudaAtomics.test_atomic_inc_global_2_32  6    11"))<
C)sJQ7JKr   c                     | j                  t        j                        \  }}d}| j                  |||ddt               y r  )r  r  r   r  r   r  s       r   test_atomic_inc_global_2_64z+TestCudaAtomics.test_atomic_inc_global_2_64  r   r   c                    |j                         } t        j                  |      |      }	 |	||f   |||       t        j                  j                  |t        j                  |dk(  |t        j                  ||kD  ||dz
                     y r2   r  r  s
             r   check_dec_indexzTestCudaAtomics.check_dec_index      xxz!DHHSM$'	#	'7"#Cf5


RXXdai.0hhtf}7=7;ax/A&B 	Cr   c                    |j                         } t        j                  |      |      }	 |	||f   |||       t        j                  j                  |t        j                  |dk(  |t        j                  ||kD  ||dz
                     y r2   r  r  s
             r   check_dec_index2z TestCudaAtomics.check_dec_index2  r  r   c                    |j                         } t        j                  |      |      } |||f   ||       t        j                  j                  |t        j                  |dk(  |t        j                  ||kD  ||dz
                     y r2   r  r  s	            r   	check_deczTestCudaAtomics.check_dec  s}    xxz!DHHSM$'	#	'7"#C0


RXXdai.0hhtf}7=7;ax/A&B 	Cr   c           	          | j                  t        j                        \  }}}d}| j                  ||||ddt               y r  )r  r  r   r  r   r  s        r   test_atomic_dec_32z"TestCudaAtomics.test_atomic_dec_32  r  r   c           	          | j                  t        j                        \  }}}d}| j                  ||||ddt               y r  )r  r  r   r  r   r  s        r   test_atomic_dec_64z"TestCudaAtomics.test_atomic_dec_64  r  r   c                     | j                  t        j                        \  }}d}| j                  |||ddt               y r  )r  r  r   r	  r   r  s       r   test_atomic_dec2_32z#TestCudaAtomics.test_atomic_dec2_32  r  r   c                     | j                  t        j                        \  }}d}| j                  |||ddt               y r  )r  r  r   r	  r   r  s       r   test_atomic_dec2_64z#TestCudaAtomics.test_atomic_dec2_64  r  r   c                     | j                  t        j                        \  }}d}| j                  |||ddt               y r  )r  r  r   r	  r   r  s       r   test_atomic_dec3_newz$TestCudaAtomics.test_atomic_dec3_new  r  r   c           	          | j                  t        j                        \  }}}d}| j                  ||||ddt               y r  )r  r  r   r  r   r  s        r   test_atomic_dec_global_32z)TestCudaAtomics.test_atomic_dec_global_32!  r  r   c           	          | j                  t        j                        \  }}}d}| j                  ||||ddt               y r  )r  r  r   r  r   r  s        r   test_atomic_dec_global_64z)TestCudaAtomics.test_atomic_dec_global_64'  r  r   c                     | j                  t        j                        \  }}d}| j                  |||ddt               y r  )r  r  r   r	  r   r  s       r   test_atomic_dec_global2_32z*TestCudaAtomics.test_atomic_dec_global2_32-  r   r   c                     | j                  t        j                        \  }}d}| j                  |||ddt               y r  )r  r  r   r	  r   r  s       r   test_atomic_dec_global2_64z*TestCudaAtomics.test_atomic_dec_global2_642  r   r   c                    t         j                  j                  ddt         j                        }t         j                  j                  ddd      j	                  t         j                        }t        j
                  dt         j                        } t        j                  d      t              } |d   |||       t         j                  j                  ||       y )	N2   d   r%  r   rA   r#  r  r$  )r  r  r'  r   r(  r  r   r*  r   rc  rd  )r  r  r"   r#   ro  s        r   test_atomic_exchz TestCudaAtomics.test_atomic_exch7  s    YY&&r3bii&@
ii2B/66ryyAii")),BDHHAB;O		%c:.


Z0r   c                    t         j                  j                  ddt         j                        }t         j                  j                  ddd      j	                  t         j                        j                  dd      } t        j                  d	      t              } |d
   ||       t         j                  j                  ||       y )Nr  r  r%  r   rA   r#  rM   rN   r  r8  )r  r  r'  r   r(  r9  r   r*  r   rc  rd  r  r  r"   ro  s       r   test_atomic_exch2z!TestCudaAtomics.test_atomic_exch2A      YY&&r3bii&@
ii2B/66ryyAII!QO9DHH89,G		)S*-


Z0r   c                    t         j                  j                  ddt         j                        }t         j                  j                  ddd      j	                  t         j                        j                  dd      } t        j                  d	      t              } |d
   ||       t         j                  j                  ||       y )Nr  r  r%  r   rA   r#  rM   rN   r  r8  )r  r  r'  r   r(  r9  r   r*  r   rc  rd  r!  s       r   test_atomic_exch3z!TestCudaAtomics.test_atomic_exch3I  r#  r   c                    t         j                  j                  ddt         j                        }t        j                  dt         j                        }t         j                  j                  dddt         j                        }d} t        j                  |      t              } |d   |||       t         j                  j                  ||       y )	Nr  r  r%  rA   r   r`  r  r$  )
r  r  r'  r   r  r   r*  r   rc  rd  )r  r  r#   r"   rs  ro  s         r   test_atomic_exch_globalz'TestCudaAtomics.test_atomic_exch_globalQ  s    YY&&r3bii&@
ii")),ii2Bbii@2!DHHSM"45		%c:.


Z0r   c                 \   t         j                  j                  ||d      j                  |      }t        j                  d|j
                        }t        j                  t              } |d   ||       t        j                  |      }t         j                  j                  ||       y )NrA   rA   r#  r3   r%  )r  r  r'  r(  r+  r&  r   r*  
atomic_maxmaxrc  rd  r  r&  lohivalsr  ro  r3  s           r   check_atomic_maxz TestCudaAtomics.check_atomic_max[  s|    yy  Rh 7>>uEhhq

+HHZ(		&#t$vvd|


T*r   c                 H    | j                  t        j                  dd       y N   r&  r-  r.  )r0  r  r  r  s    r   test_atomic_max_int32z%TestCudaAtomics.test_atomic_max_int32c      BHHEBr   c                 H    | j                  t        j                  dd       y Nr   r4  r5  )r0  r  r   r6  s    r   test_atomic_max_uint32z&TestCudaAtomics.test_atomic_max_uint32f      BII!>r   c                 H    | j                  t        j                  dd       y r2  )r0  r  rb  r6  s    r   test_atomic_max_int64z%TestCudaAtomics.test_atomic_max_int64i  r8  r   c                 H    | j                  t        j                  dd       y r:  )r0  r  r   r6  s    r   test_atomic_max_uint64z&TestCudaAtomics.test_atomic_max_uint64l  r<  r   c                 H    | j                  t        j                  dd       y r2  )r0  r  r   r6  s    r   test_atomic_max_float32z'TestCudaAtomics.test_atomic_max_float32o      BJJ6eDr   c                 H    | j                  t        j                  dd       y r2  )r0  r  r   r6  s    r   test_atomic_max_doublez&TestCudaAtomics.test_atomic_max_doubler  rC  r   c                    t         j                  j                  ddd      j                  t         j                        }t        j
                  dt         j                        } t        j                  d      t              } |d   ||       t        j                  |      }t         j                  j                  ||       y Nr   r4  r)  r#  r3   void(float64[:], float64[:,:]))r  r  r'  r(  r   r+  r   r*  !atomic_max_double_normalizedindexr+  rc  rd  r  r/  r  ro  r3  s        r   &test_atomic_max_double_normalizedindexz6TestCudaAtomics.test_atomic_max_double_normalizedindexu  s    yy  E 9@@Lhhq"**%>DHH=>-/		&#t$vvd|


T*r   c                    t         j                  j                  ddd      j                  t         j                        }t        j
                  dt         j                        } t        j                  d      t              } |d   ||       t        j                  |      }t         j                  j                  ||       y Nr      rA   r#  r3   void(float64[:], float64[:])r$  )r  r  r'  r(  r   r+  r   r*  atomic_max_double_oneindexr+  rc  rd  rJ  s        r   test_atomic_max_double_oneindexz/TestCudaAtomics.test_atomic_max_double_oneindex  s    yy  Cb 188Dhhq"**%<DHH;<&(		%d#vvd|


T*r   c                 ^   t         j                  j                  ||d      j                  |      }t        j                  dg|j
                        }t        j                  t              } |d   ||       t        j                  |      }t         j                  j                  ||       y )Nr)  r#  r4  r%  )r  r  r'  r(  r    r&  r   r*  
atomic_minminrc  rd  r,  s           r   check_atomic_minz TestCudaAtomics.check_atomic_min  s~    yy  Rh 7>>uEhhwdjj1HHZ(		&#t$vvd|


T*r   c                 H    | j                  t        j                  dd       y r2  )rU  r  r  r6  s    r   test_atomic_min_int32z%TestCudaAtomics.test_atomic_min_int32  r8  r   c                 H    | j                  t        j                  dd       y r:  )rU  r  r   r6  s    r   test_atomic_min_uint32z&TestCudaAtomics.test_atomic_min_uint32  r<  r   c                 H    | j                  t        j                  dd       y r2  )rU  r  rb  r6  s    r   test_atomic_min_int64z%TestCudaAtomics.test_atomic_min_int64  r8  r   c                 H    | j                  t        j                  dd       y r:  )rU  r  r   r6  s    r   test_atomic_min_uint64z&TestCudaAtomics.test_atomic_min_uint64  r<  r   c                 H    | j                  t        j                  dd       y r2  )rU  r  r   r6  s    r   test_atomic_min_floatz%TestCudaAtomics.test_atomic_min_float  rC  r   c                 H    | j                  t        j                  dd       y r2  )rU  r  r   r6  s    r   test_atomic_min_doublez&TestCudaAtomics.test_atomic_min_double  rC  r   c                    t         j                  j                  ddd      j                  t         j                        }t        j
                  dt         j                        dz  } t        j                  d      t              } |d   ||       t        j                  |      }t         j                  j                  ||       y rG  )r  r  r'  r(  r   onesr   r*  !atomic_min_double_normalizedindexrT  rc  rd  rJ  s        r   &test_atomic_min_double_normalizedindexz6TestCudaAtomics.test_atomic_min_double_normalizedindex  s    yy  E 9@@Lgga$u,>DHH=>-/		&#t$vvd|


T*r   c                    t         j                  j                  ddd      j                  t         j                        }t        j
                  dt         j                        dz  } t        j                  d      t              } |d   ||       t        j                  |      }t         j                  j                  ||       y rM  )r  r  r'  r(  r   rc  r   r*  atomic_min_double_oneindexrT  rc  rd  rJ  s        r   test_atomic_min_double_oneindexz/TestCudaAtomics.test_atomic_min_double_oneindex  s    yy  Cb 188Dgga$s*<DHH;<&(		%d#vvd|


T*r   c                     t        j                  d      |      }t        j                  j	                  ddd      j                  t        j                        }t        j                  dt        j                        t        j                  z   } |d   ||       t        j                  j                  |t        j                  g       y )NrH  r   rN  r3   r3   r#  r3   )r   r*  r  r  r'  r(  r   r+  nanrc  rd  )r  r   ro  r/  r  s        r    _test_atomic_minmax_nan_locationz0TestCudaAtomics._test_atomic_minmax_nan_location  s    >DHH=>tD	yy  Ce 4;;BJJGhhq"**%.	$T"


bffX.r   c                     t        j                  d      |      }t        j                  j	                  ddd      j                  t        j                        }|j                         }t        j                  dt        j                        t        j                  z   } |d   ||       t        j                  j                  ||       y )NrH  r   rN  r3   r#  rj  )r   r*  r  r  r'  r(  r   r)  r+  rk  rc  rd  )r  r   ro  r  r3  r/  s         r   _test_atomic_minmax_nan_valz+TestCudaAtomics._test_atomic_minmax_nan_val  s    >DHH=>tD	ii3Q/66rzzBxxzxx

+bff4	$T"


T*r   c                 .    | j                  t               y r   )rl  rS  r6  s    r   test_atomic_min_nan_locationz,TestCudaAtomics.test_atomic_min_nan_location      --j9r   c                 .    | j                  t               y r   )rl  r*  r6  s    r   test_atomic_max_nan_locationz,TestCudaAtomics.test_atomic_max_nan_location  rq  r   c                 .    | j                  t               y r   )rn  rS  r6  s    r   test_atomic_min_nan_valz'TestCudaAtomics.test_atomic_min_nan_val      ((4r   c                 .    | j                  t               y r   )rn  r*  r6  s    r   test_atomic_max_nan_valz'TestCudaAtomics.test_atomic_max_nan_val  rv  r   c                    t         j                  j                  ddd      j                  t         j                        }t        j
                  dt         j                        }d} t        j                  |      t              } |d   ||       t        j                  |      }t         j                  j                  ||       y Nr   rA   r#  r3   rO  r$  )r  r  r'  r(  r   r+  r   r*  atomic_max_double_sharedr+  rc  rd  r  r/  r  rs  ro  r3  s         r   test_atomic_max_double_sharedz-TestCudaAtomics.test_atomic_max_double_shared  s    yy  BR 077

Chhq"**%,!DHHSM":;		%d#vvd|


T*r   c                    t         j                  j                  ddd      j                  t         j                        }t        j
                  dt         j                        dz  }d} t        j                  |      t              } |d   ||       t        j                  |      }t         j                  j                  ||       y rz  )r  r  r'  r(  r   rc  r   r*  atomic_min_double_sharedrT  rc  rd  r|  s         r   test_atomic_min_double_sharedz-TestCudaAtomics.test_atomic_min_double_shared  s    yy  BR 077

Cgga$r),!DHHSM":;		%d#vvd|


T*r   c                    |g|dz  z  |g|dz  z  z   }t         j                  j                  |       t        j                  ||      }|dk(  rd|_        t        j
                  |      }t         j                  j                  dd|j                        j                  |j                        }	||k(  }
||k(  }t        j
                  |      }|	|
   ||
<   |||<   |j                         }t        j                  |      }|dk(  r |d   |||	|       n |d   |||	|       t         j                  j                  ||       t         j                  j                  ||       y )	Nr  r%  )
   r3   r  r#  r  r  )r  r  )r  r  shuffleasarrayr;   
zeros_liker'  r(  r&  r)  r   r*  rc  assert_array_equal)r  nfillunfillr&  cas_funcndimr  outr"   	fill_maskunfill_mask
expect_res
expect_outro  s                  r   	check_caszTestCudaAtomics.check_cas  s9   fQ6(a1f"55
		#jjE*19 CImmC ii2CII6==ciiH4K	Vm]]3'
 #I
9"(
;XXZ
HHX&	19Ifc3T2)I()#sC>


%%j#6


%%j#6r   c                 T    | j                  dddt        j                  t               y Nr  r  r  r  r  r&  r  )r  r  r  r  r6  s    r   test_atomic_compare_and_swapz,TestCudaAtomics.test_atomic_compare_and_swap  "    3r 7 	 	9r   c                 T    | j                  dddt        j                  t               y Nr  r  r  )r  r  rb  r  r6  s    r   test_atomic_compare_and_swap2z-TestCudaAtomics.test_atomic_compare_and_swap2  r  r   c                    t         j                  j                  ddt         j                        }t         j                  j                  ddt         j                        }| j	                  d||t         j                  t
               y Nr  r  r%  r3      r  r  )r  r  r'  r   r  r  r  rfillrunfills      r   test_atomic_compare_and_swap3z-TestCudaAtomics.test_atomic_compare_and_swap3  `    		!!"c!;))##Ar#;5		 7 	 	9r   c                    t         j                  j                  ddt         j                        }t         j                  j                  ddt         j                        }| j	                  d||t         j                  t
               y r  )r  r  r'  r   r  r  r  s      r   test_atomic_compare_and_swap4z-TestCudaAtomics.test_atomic_compare_and_swap4  r  r   c                 T    | j                  dddt        j                  t               y r  )r  r  r  r  r6  s    r   test_atomic_cas_1dimz$TestCudaAtomics.test_atomic_cas_1dim%  "    3r / 	 	1r   c                 V    | j                  dddt        j                  t        d       y )Nr  r  r  r  r  r  r  r&  r  r  )r  r  r  r  r6  s    r   test_atomic_cas_2dimz$TestCudaAtomics.test_atomic_cas_2dim)  $    3r /a 	 	9r   c                 T    | j                  dddt        j                  t               y r  )r  r  rb  r  r6  s    r   test_atomic_cas2_1dimz%TestCudaAtomics.test_atomic_cas2_1dim-  r  r   c                 V    | j                  dddt        j                  t        d       y )Nr  r  r  r  r  )r  r  rb  r  r6  s    r   test_atomic_cas2_2dimz%TestCudaAtomics.test_atomic_cas2_2dim1  r  r   c                    t         j                  j                  ddt         j                        }t         j                  j                  ddt         j                        }| j	                  d||t         j                  t
               y r  )r  r  r'  r   r  r  r  s      r   test_atomic_cas3_1dimz%TestCudaAtomics.test_atomic_cas3_1dim5  `    		!!"c!;))##Ar#;5		 / 	 	1r   c                    t         j                  j                  ddt         j                        }t         j                  j                  ddt         j                        }| j	                  d||t         j                  t
        d       y 	Nr  r  r%  r3   r  r  r  r  )r  r  r'  r   r  r  r  s      r   test_atomic_cas3_2dimz%TestCudaAtomics.test_atomic_cas3_2dim;  b    		!!"c!;))##Ar#;5		 /a 	 	9r   c                    t         j                  j                  ddt         j                        }t         j                  j                  ddt         j                        }| j	                  d||t         j                  t
               y r  )r  r  r'  r   r  r  r  s      r   test_atomic_cas4_1dimz%TestCudaAtomics.test_atomic_cas4_1dimA  r  r   c                    t         j                  j                  ddt         j                        }t         j                  j                  ddt         j                        }| j	                  d||t         j                  t
        d       y r  )r  r  r'  r   r  r  r  s      r   test_atomic_cas4_2dimz%TestCudaAtomics.test_atomic_cas4_2dimG  r  r   c                    t        j                  dt         j                        }||d<    |d   |       t        j                  |      r(| j	                  t        j                  |d                y | j                  |d   |       y )Nr  r%  r   rj  r3   )r  r+  r   isnanr-  assertEqualr  rZ  initialr   s       r   _test_atomic_returns_oldz(TestCudaAtomics._test_atomic_returns_oldR  sd    HHQbjj)!tQ88GOOBHHQqTN+QqT7+r   c                 T    t         j                  d        }| j                  |d       y )Nc                 L    t         j                  j                  | dd      | d<   y r2   )r   rC   rD   r   s    r   rZ  z;TestCudaAtomics.test_atomic_add_returns_old.<locals>.kernel\      ;;??1a+AaDr   r  r   r*  r  r  rZ  s     r   test_atomic_add_returns_oldz+TestCudaAtomics.test_atomic_add_returns_old[  *    		, 
	, 	%%fb1r   c                 T    t         j                  d        }| j                  |d       y )Nc                 L    t         j                  j                  | dd      | d<   y r2   r   rC   r+  r  s    r   rZ  zBTestCudaAtomics.test_atomic_max_returns_no_replace.<locals>.kernelc  r  r   r  r  r  s     r   "test_atomic_max_returns_no_replacez2TestCudaAtomics.test_atomic_max_returns_no_replaceb  r  r   c                 T    t         j                  d        }| j                  |d       y )Nc                 L    t         j                  j                  | dd      | d<   y Nr   r  r3   r  r  s    r   rZ  zCTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernelj      ;;??1a,AaDr   r3   r  r  s     r   #test_atomic_max_returns_old_replacez3TestCudaAtomics.test_atomic_max_returns_old_replacei  s*    		- 
	- 	%%fa0r   c                 p    t         j                  d        }| j                  |t        j                         y )Nc                 L    t         j                  j                  | dd      | d<   y r2   r  r  s    r   rZ  zHTestCudaAtomics.test_atomic_max_returns_old_nan_in_array.<locals>.kernelq  r  r   r   r*  r  r  rk  r  s     r   (test_atomic_max_returns_old_nan_in_arrayz8TestCudaAtomics.test_atomic_max_returns_old_nan_in_arrayp  s.    		, 
	, 	%%fbff5r   c                 T    t         j                  d        }| j                  |d       y )Nc                 h    t         j                  j                  | dt        j                        | d<   y r2   )r   rC   r+  r  rk  r  s    r   rZ  zCTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernelx       ;;??1a0AaDr   r  r  r  s     r   #test_atomic_max_returns_old_nan_valz3TestCudaAtomics.test_atomic_max_returns_old_nan_valw  *    		1 
	1 	%%fb1r   c                 T    t         j                  d        }| j                  |d       y )Nc                 L    t         j                  j                  | dd      | d<   y Nr      r3   r   rC   rT  r  s    r   rZ  zFTestCudaAtomics.test_atomic_min_returns_old_no_replace.<locals>.kernel  r  r   r  r  r  s     r   &test_atomic_min_returns_old_no_replacez6TestCudaAtomics.test_atomic_min_returns_old_no_replace~  *    		- 
	- 	%%fb1r   c                 T    t         j                  d        }| j                  |d       y )Nc                 L    t         j                  j                  | dd      | d<   y r  r  r  s    r   rZ  zCTestCudaAtomics.test_atomic_min_returns_old_replace.<locals>.kernel  r  r   r  r  r  s     r   #test_atomic_min_returns_old_replacez3TestCudaAtomics.test_atomic_min_returns_old_replace  r  r   c                 p    t         j                  d        }| j                  |t        j                         y )Nc                 L    t         j                  j                  | dd      | d<   y r  r  r  s    r   rZ  zHTestCudaAtomics.test_atomic_min_returns_old_nan_in_array.<locals>.kernel  r  r   r  r  s     r   (test_atomic_min_returns_old_nan_in_arrayz8TestCudaAtomics.test_atomic_min_returns_old_nan_in_array  s.    		- 
	- 	%%fbff5r   c                 T    t         j                  d        }| j                  |d       y )Nc                 h    t         j                  j                  | dt        j                        | d<   y r2   )r   rC   rT  r  rk  r  s    r   rZ  zCTestCudaAtomics.test_atomic_min_returns_old_nan_val.<locals>.kernel  r  r   r  r  r  s     r   #test_atomic_min_returns_old_nan_valz3TestCudaAtomics.test_atomic_min_returns_old_nan_val  r  r   c                 l   t         j                  j                  ||d      j                  |      }||dd d<   t        j                  d|j
                        }t        j                  t              } |d   ||       t        j                  |      }t         j                  j                  ||       y )Nr)  r#  r3   r  r%  )r  r  r'  r(  r+  r&  r   r*  atomic_nanmaxnanmaxrc  rd  	r  r&  r-  r.  init_valr/  r  ro  r3  s	            r   check_atomic_nanmaxz#TestCudaAtomics.check_atomic_nanmax  s    yy  Rh 7>>uEQTT
hhq

+HH]+		&#t$yy


T*r   c                 J    | j                  t        j                  ddd       y Nr3  r4  r   r&  r-  r.  r  )r  r  r  r6  s    r   test_atomic_nanmax_int32z(TestCudaAtomics.test_atomic_nanmax_int32  "      rxxFu*+ 	! 	-r   c                 J    | j                  t        j                  ddd       y Nr   r4  r  )r  r  r   r6  s    r   test_atomic_nanmax_uint32z)TestCudaAtomics.test_atomic_nanmax_uint32  "      ryyQ5*+ 	! 	-r   c                 J    | j                  t        j                  ddd       y r  )r  r  rb  r6  s    r   test_atomic_nanmax_int64z(TestCudaAtomics.test_atomic_nanmax_int64  r  r   c                 J    | j                  t        j                  ddd       y r  )r  r  r   r6  s    r   test_atomic_nanmax_uint64z)TestCudaAtomics.test_atomic_nanmax_uint64  r  r   c                 f    | j                  t        j                  ddt        j                         y Nr3  r4  r  )r  r  r   rk  r6  s    r   test_atomic_nanmax_float32z*TestCudaAtomics.test_atomic_nanmax_float32  &      rzzf*,&& 	! 	2r   c                 f    | j                  t        j                  ddt        j                         y r  )r  r  r   rk  r6  s    r   test_atomic_nanmax_doublez)TestCudaAtomics.test_atomic_nanmax_double  r  r   c                    t         j                  j                  ddd      j                  t         j                        }t         j
                  |dd d<   t        j                  dg|j                        }d} t        j                  |      t              } |d   ||       t        j                  |      }t         j                  j                  ||       y 	Nr   rA   r#  r3   r  r%  rO  r$  )r  r  r'  r(  r   rk  r    r&  r   r*  atomic_nanmax_double_sharedr  rc  rd  r|  s         r    test_atomic_nanmax_double_sharedz0TestCudaAtomics.test_atomic_nanmax_double_shared  s    yy  BR 077

CVVQTT
hhs$**-,!DHHSM"=>		%d#yy


T*r   c                    t         j                  j                  ddd      j                  t         j                        }t         j
                  |dd d<   t        j                  dt         j                        } t        j                  d      t              } |d   ||       t        j                  |      }t         j                  j                  ||       y 	Nr   rN  rA   r#  r3   r  rO  r$  )r  r  r'  r(  r   rk  r+  r   r*  rP  r  rc  rd  rJ  s        r   "test_atomic_nanmax_double_oneindexz2TestCudaAtomics.test_atomic_nanmax_double_oneindex  s    yy  Cb 188DVVQTT
hhq"**%<DHH;<&(		%d#yy


T*r   c                 n   t         j                  j                  ||d      j                  |      }||dd d<   t        j                  dg|j
                        }t        j                  t              } |d   ||       t        j                  |      }t         j                  j                  ||       y )Nr)  r#  r3   r  r4  r%  )r  r  r'  r(  r    r&  r   r*  atomic_nanminnanminrc  rd  r  s	            r   check_atomic_nanminz#TestCudaAtomics.check_atomic_nanmin  s    yy  Rh 7>>uEQTT
hhwdjj1HH]+		&#t$yy


T*r   c                 J    | j                  t        j                  ddd       y r  )r  r  r  r6  s    r   test_atomic_nanmin_int32z(TestCudaAtomics.test_atomic_nanmin_int32  r  r   c                 J    | j                  t        j                  ddd       y r  )r  r  r   r6  s    r   test_atomic_nanmin_uint32z)TestCudaAtomics.test_atomic_nanmin_uint32  r  r   c                 J    | j                  t        j                  ddd       y r  )r  r  rb  r6  s    r   test_atomic_nanmin_int64z(TestCudaAtomics.test_atomic_nanmin_int64  r  r   c                 J    | j                  t        j                  ddd       y r  )r  r  r   r6  s    r   test_atomic_nanmin_uint64z)TestCudaAtomics.test_atomic_nanmin_uint64  r  r   c                 f    | j                  t        j                  ddt        j                         y r  )r  r  r   rk  r6  s    r   test_atomic_nanmin_floatz(TestCudaAtomics.test_atomic_nanmin_float  r  r   c                 f    | j                  t        j                  ddt        j                         y r  )r  r  r   rk  r6  s    r   test_atomic_nanmin_doublez)TestCudaAtomics.test_atomic_nanmin_double  r  r   c                    t         j                  j                  ddd      j                  t         j                        }t         j
                  |dd d<   t        j                  dg|j                        }d} t        j                  |      t              } |d   ||       t        j                  |      }t         j                  j                  ||       y r  )r  r  r'  r(  r   rk  r    r&  r   r*  atomic_nanmin_double_sharedr  rc  rd  r|  s         r    test_atomic_nanmin_double_sharedz0TestCudaAtomics.test_atomic_nanmin_double_shared  s    yy  BR 077

CVVQTT
hht4::.,!DHHSM"=>		%d#yy


T*r   c                    t         j                  j                  ddd      j                  t         j                        }t         j
                  |dd d<   t        j                  dgt         j                        } t        j                  d      t              } |d   ||       t        j                  |      }t         j                  j                  ||       y r  )r  r  r'  r(  r   rk  r    r   r*  rg  r  rc  rd  rJ  s        r   "test_atomic_nanmin_double_oneindexz2TestCudaAtomics.test_atomic_nanmin_double_oneindex  s    yy  Cb 188DVVQTT
hhubjj)<DHH;<&(		%d#yy


T*r   c                    t        j                  dt         j                        }||d<   t         j                  |d<    |d   |       t        j                  |      rO| j                  t        j                  |d                | j                  t        j                  |d                y | j                  |d   |       y )Nr  r%  r   r3   rj  )r  r+  r   rk  r  assertFalser-  r  r  s       r   _test_atomic_nan_returns_oldz,TestCudaAtomics._test_atomic_nan_returns_old  s    HHQbjj)!vv!tQ88GRXXad^,OOBHHQqTN+QqT7+r   c                 T    t         j                  d        }| j                  |d       y )Nc                 L    t         j                  j                  | dd      | d<   y r2   r   rC   r  r  s    r   rZ  zITestCudaAtomics.test_atomic_nanmax_returns_old_no_replace.<locals>.kernel      ;;%%aA.AaDr   r  r   r*  r%  r  s     r   )test_atomic_nanmax_returns_old_no_replacez9TestCudaAtomics.test_atomic_nanmax_returns_old_no_replace  s*    		/ 
	/ 	))&"5r   c                 T    t         j                  d        }| j                  |d       y )Nc                 L    t         j                  j                  | dd      | d<   y r  r(  r  s    r   rZ  zFTestCudaAtomics.test_atomic_nanmax_returns_old_replace.<locals>.kernel"      ;;%%aB/AaDr   r3   r*  r  s     r   &test_atomic_nanmax_returns_old_replacez6TestCudaAtomics.test_atomic_nanmax_returns_old_replace!  s*    		0 
	0 	))&!4r   c                 p    t         j                  d        }| j                  |t        j                         y )Nc                 L    t         j                  j                  | dd      | d<   y r2   r(  r  s    r   rZ  zKTestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array.<locals>.kernel)  r)  r   r   r*  r%  r  rk  r  s     r   +test_atomic_nanmax_returns_old_nan_in_arrayz;TestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array(  s.    		/ 
	/ 	))&"&&9r   c                 T    t         j                  d        }| j                  |d       y )Nc                 h    t         j                  j                  | dt        j                        | d<   y r2   )r   rC   r  r  rk  r  s    r   rZ  zFTestCudaAtomics.test_atomic_nanmax_returns_old_nan_val.<locals>.kernel0  "    ;;%%aBFF3AaDr   r  r*  r  s     r   &test_atomic_nanmax_returns_old_nan_valz6TestCudaAtomics.test_atomic_nanmax_returns_old_nan_val/  *    		4 
	4 	))&"5r   c                 T    t         j                  d        }| j                  |d       y )Nc                 L    t         j                  j                  | dd      | d<   y r  r   rC   r  r  s    r   rZ  zITestCudaAtomics.test_atomic_nanmin_returns_old_no_replace.<locals>.kernel7  r.  r   r  r*  r  s     r   )test_atomic_nanmin_returns_old_no_replacez9TestCudaAtomics.test_atomic_nanmin_returns_old_no_replace6  *    		0 
	0 	))&"5r   c                 T    t         j                  d        }| j                  |d       y )Nc                 L    t         j                  j                  | dd      | d<   y r  r;  r  s    r   rZ  zFTestCudaAtomics.test_atomic_nanmin_returns_old_replace.<locals>.kernel>  r.  r   r  r*  r  s     r   &test_atomic_nanmin_returns_old_replacez6TestCudaAtomics.test_atomic_nanmin_returns_old_replace=  r=  r   c                 p    t         j                  d        }| j                  |t        j                         y )Nc                 L    t         j                  j                  | dd      | d<   y r  r;  r  s    r   rZ  zKTestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_array.<locals>.kernelE  r.  r   r2  r  s     r   +test_atomic_nanmin_returns_old_nan_in_arrayz;TestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_arrayD  s.    		0 
	0 	))&"&&9r   c                 T    t         j                  d        }| j                  |d       y )Nc                 h    t         j                  j                  | dt        j                        | d<   y r2   )r   rC   r  r  rk  r  s    r   rZ  zFTestCudaAtomics.test_atomic_nanmin_returns_old_nan_val.<locals>.kernelL  r6  r   r  r*  r  s     r   &test_atomic_nanmin_returns_old_nan_valz6TestCudaAtomics.test_atomic_nanmin_returns_old_nan_valK  r8  r   )T)r3   )__name__
__module____qualname__r  r5  r<  r@  rF  rJ  rL  r]  rg  rl  rp  ru  rx  rz  r}  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r	  r  r  r  r  r  r  r  r  r  r  r"  r%  r'  r0  r7  r;  r>  r@  rB  rE  rK  rQ  rU  rW  rY  r[  r]  r_  ra  re  rh  rl  rn  rp  rs  ru  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  r   r  r  r
  r  r  r  r  r  r  r  r  r   r"  r%  r+  r/  r3  r7  r<  r@  rC  rF  __classcell__)r   s   @r   r  r    s   2$612$61322(7 4G*G"B
-11
-11
+//+//-::+8-::+8-::+8%
LLLM
M
G
G
D
11L
L
CCCM
M
G
G
D
11L
L
1111+C?C?EE+++C?C?EE++*/+::55++76999919191919,221622262+----22	+	++----22	+	+	,65:666:6r   r  __main__)qnumpyr  textwrapr   numbar   r   r   r   r   numba.cuda.testingr	   r
   r   
numba.corer   r*  r   r   r   r.   r0   r9   r<   r>   rF   rI   rP   rS   rU   r\   r`   rd   rh   rk   rq   rt   rx   rz   r|   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  r*  rI  rP  r{  rS  rd  rg  r  r  $atomic_nanmax_double_normalizedindexatomic_nanmax_double_oneindexr	  r  $atomic_nanmin_double_normalizedindexatomic_nanmin_double_oneindexr  r  r  r  r  rG  mainr   r   r   <module>rV     s=     8 8 D D  
  
  
  
  
	 	 
  
  
 K
J
H
G
M
O
N
H
G
M
IHPO%
M
L
H
G
M
K
H
M
O
H
M
M
H
M
I%
M
L
I
N
J7
K
H
M
I7
K
H
M
I7
B
D
H
H
M
I7
B
D
H
H
M
I7
C
I
I
J!FJ 66GH.0J56GH.0J 12=4 ; 12=4 ;OAA{6l {6|# zHMMO r   