
    J/PhQ                     ^   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jp                     dS dS )\    N)dedent)cudauint32uint64float32float64)unittestCUDATestCasecc_X_or_above)configT)devicec                      t          |           S N)r   nums    d/var/www/html/test/jupyter/venv/lib/python3.11/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    s88Or   c                     | S r    r   s    r   atomic_cast_noner      s    Jr   c	                 $   t           j        j        }	t           j                            ||          }
||
|	<   t          j                      |||	         |z            }|r||z  } ||
||           t          j                     |
|	         | |	<   d S r   r   	threadIdxxsharedarraysyncthreads)aryidxop2	ary_dtypeary_nelements
binop_func	cast_funcinitializerneg_idxtidsmbins               r   atomic_binary_1dim_sharedr.      s     .
C			=)	4	4BBsG
)CH},
-
-C "M!Jr3#wCHHHr   c                 "   t           j        j        }t           j                            ||          }| |         ||<   t          j                      |||         |z            }	 |||	|           t          j                     ||         | |<   d S r   r   )
r"   r#   r$   r%   r&   r'   r(   r+   r,   r-   s
             r   atomic_binary_1dim_shared2r0   (   s     .
C			=)	4	4B#hBsG
)CH},
-
-CJr3#wCHHHr   c                    t           j        j        }t           j        j        }t           j                            ||          }	| ||f         |	||f<   t          j                     | ||          f}
|r"|
d         |d         z  |
d         |d         z  f}
 ||	|
|           t          j                     |	||f         | ||f<   d S 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	0	0BRVBr2vJ{{2
C =1v	!$c!fy|&;<Jr3RV*CBKKKr   c                     t           j        j        }t           j        j        }| ||          f}|r,|d         | j        d         z  |d         | j        d         z  f} || ||           d S r2   )r   r   r   r4   shape)r"   r$   r'   r6   r*   r7   r8   r-   s           r   atomic_binary_2dim_globalr<   E   sr    		B		B{{2
C =1v	!$c!fsy|&;<JsCr   c                     t           j        j        }t          ||         |z            }|r||z  } || ||           d S r   )r   r   r   r   )r"   r#   r&   r$   r'   r*   r+   r-   s           r   atomic_binary_1dim_globalr>   O   sP     .
C
c#h&
'
'C "M!JsCr   c                 f    t          | | dt          dt          j        j        t
          dd	  	         d S Nr3       r   Fr.   r   r   atomicaddr   r"   s    r   
atomic_addrF   Y   ;    c362"ko/?EK K K K Kr   c                 f    t          | | dt          dt          j        j        t
          dd	  	         d S )Nr3   rA   r   TrB   rE   s    r   atomic_add_wraprI   ^   s;    c362"ko/?DJ J J J Jr   c           	      b    t          | dt          dt          j        j        t
          d           d S Nr3         Fr9   r   r   rC   rD   r   rE   s    r   atomic_add2rP   c   7    c1ff"ko/?H H H H Hr   c           	      b    t          | dt          dt          j        j        t
          d           d S )Nr3   rL   TrO   rE   s    r   atomic_add2_wraprS   h   s7    c1ff"ko/?G G G G Gr   c           	      b    t          | dt          dt          j        j        t
          d           d S rK   )r9   r   r   rC   rD   r   rE   s    r   atomic_add3rU   m   7    c1ff"ko/DeM M M M Mr   c                 f    t          | | dt          dt          j        j        t
          dd	  	         d S N      ?rA           Fr.   r   r   rC   rD   r   rE   s    r   atomic_add_floatr\   r   ;    c3Wb"ko/A3O O O O Or   c                 f    t          | | dt          dt          j        j        t
          dd	  	         d S NrY   rA   rZ   Tr[   rE   s    r   atomic_add_float_wrapr`   w   s;    c3Wb"ko/A3N N N N Nr   c           	      b    t          | dt          dt          j        j        t
          d           d S NrY   rL   Fr9   r   r   rC   rD   r   rE   s    r   atomic_add_float_2rd   |   7    c3"ko/?H H H H Hr   c           	      b    t          | dt          dt          j        j        t
          d           d S NrY   rL   Trc   rE   s    r   atomic_add_float_2_wraprh      7    c3"ko/?G G G G Gr   c           	      b    t          | dt          dt          j        j        t
          d           d S rb   )r9   r   r   rC   rD   r   rE   s    r   atomic_add_float_3rk      7    c3"ko/DeM M M M Mr   c                 L    t          || ddt          j        j        d           d S NrA   rY   Fr>   r   rC   rD   r#   r"   s     r   atomic_add_double_globalrq      $    c3C%HHHHHr   c                 L    t          || ddt          j        j        d           d S )NrA   rY   Tro   rp   s     r   atomic_add_double_global_wraprt      s$    c3C$GGGGGr   c                 T    t          | dt          j        j        t          d           d S Nr3   Fr<   r   rC   rD   r   rE   s    r   atomic_add_double_global_2rx      s#    c1dko7GOOOOOr   c                 T    t          | dt          j        j        t          d           d S )Nr3   Trw   rE   s    r   atomic_add_double_global_2_wraprz      s#    c1dko7GNNNNNr   c                 T    t          | dt          j        j        t          d           d S rv   )r<   r   rC   rD   r   rE   s    r   atomic_add_double_global_3r|      s.    c1dko7L#% % % % %r   c                 f    t          || dt          dt          j        j        t
          dd	  	         d S rX   r.   r   r   rC   rD   r   rp   s     r   atomic_add_doubler      ;    c3Wb"ko/?eM M M M Mr   c                 f    t          || dt          dt          j        j        t
          dd	  	         d S r_   r~   rp   s     r   atomic_add_double_wrapr      s;    c3Wb"ko/?dL L L L Lr   c           	      b    t          | dt          dt          j        j        t
          d           d S rb   r9   r   r   rC   rD   r   rE   s    r   atomic_add_double_2r      re   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rg   r   rE   s    r   atomic_add_double_2_wrapr      ri   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rb   )r9   r   r   rC   rD   r   rE   s    r   atomic_add_double_3r      rl   r   c                 f    t          | | dt          dt          j        j        t
          dd	  	         d S r@   )r.   r   r   rC   subr   rE   s    r   
atomic_subr      rG   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rK   )r9   r   r   rC   r   r   rE   s    r   atomic_sub2r      rQ   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rK   )r9   r   r   rC   r   r   rE   s    r   atomic_sub3r      rV   r   c                 f    t          | | dt          dt          j        j        t
          dd	  	         d S rX   )r.   r   r   rC   r   r   rE   s    r   atomic_sub_floatr      r]   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rb   )r9   r   r   rC   r   r   rE   s    r   atomic_sub_float_2r      re   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rb   )r9   r   r   rC   r   r   rE   s    r   atomic_sub_float_3r      rl   r   c                 f    t          || dt          dt          j        j        t
          dd	  	         d S rX   )r.   r   r   rC   r   r   rp   s     r   atomic_sub_doubler      r   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rb   )r9   r   r   rC   r   r   rE   s    r   atomic_sub_double_2r      re   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rb   r9   r   r   rC   r   r   rE   s    r   atomic_sub_double_3r      rl   r   c                 L    t          || ddt          j        j        d           d S rn   )r>   r   rC   r   rp   s     r   atomic_sub_double_globalr      rr   r   c                 T    t          | dt          j        j        t          d           d S )NrY   F)r<   r   rC   r   r   rE   s    r   atomic_sub_double_global_2r      s.    c39I#% % % % %r   c           	      b    t          | dt          dt          j        j        t
          d           d S rb   r   rE   s    r   atomic_sub_double_global_3r      rl   r   c                 f    t          | | |t          dt          j        j        t
          dd	  	         d S )NrA   r3   F)r.   r   r   rC   and_r   r"   r$   s     r   
atomic_andr      s<    c3VR"k.0@!UL L L L Lr   c           	      b    t          | |t          dt          j        j        t
          d           d S NrL   F)r9   r   r   rC   r   r   r   s     r   atomic_and2r      8    c3"k.0@%I I I I Ir   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   atomic_and3r      s8    c3"k.0EuN N N N Nr   c                 L    t          || d|t          j        j        d           d S NrA   F)r>   r   rC   r   r#   r"   r$   s      r   atomic_and_globalr     %    c3C1A5IIIIIr   c                 T    t          | |t          j        j        t          d           d S NF)r<   r   rC   r   r   r   s     r   atomic_and_global_2r     s.    c3(8.7 7 7 7 7r   c                 f    t          | | |t          dt          j        j        t
          dd	  	         d S NrA   r   F)r.   r   r   rC   or_r   r   s     r   	atomic_orr     ;    c3VR"ko/?EK K K K Kr   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   
atomic_or2r     7    c3"ko/?H H H H Hr   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   
atomic_or3r     7    c3"ko/DeM M M M Mr   c                 L    t          || d|t          j        j        d           d S r   )r>   r   rC   r   r   s      r   atomic_or_globalr     rr   r   c                 T    t          | |t          j        j        t          d           d S r   )r<   r   rC   r   r   r   s     r   atomic_or_global_2r     -    c3.7 7 7 7 7r   c                 f    t          | | |t          dt          j        j        t
          dd	  	         d S r   )r.   r   r   rC   xorr   r   s     r   
atomic_xorr   $  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   atomic_xor2r   )  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   atomic_xor3r   .  r   r   c                 L    t          || d|t          j        j        d           d S r   )r>   r   rC   r   r   s      r   atomic_xor_globalr   3  rr   r   c                 T    t          | |t          j        j        t          d           d S r   )r<   r   rC   r   r   r   s     r   atomic_xor_global_2r   7  r   r   c           	      b    t          | ||t          dt          j        j        t
                     d S NrA   )r0   r   r   rC   incr   r"   r#   r$   s      r   atomic_inc32r   <  7    sCfb#{0@B B B B Br   c           	      b    t          | ||t          dt          j        j        t
                     d S r   )r0   r   r   rC   r   r   r   s      r   atomic_inc64r   A  7    sCfb#{0BD D D D Dr   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   atomic_inc2_32r   F  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   atomic_inc2_64r   K  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   atomic_inc3r   P  r   r   c                 L    t          || d|t          j        j        d           d S r   )r>   r   rC   r   r   s      r   atomic_inc_globalr   U  rr   r   c                 T    t          | |t          j        j        t          d           d S r   )r<   r   rC   r   r   r   s     r   atomic_inc_global_2r   Y  r   r   c           	      b    t          | ||t          dt          j        j        t
                     d S r   )r0   r   r   rC   decr   r   s      r   atomic_dec32r   ^  r   r   c           	      b    t          | ||t          dt          j        j        t
                     d S r   )r0   r   r   rC   r   r   r   s      r   atomic_dec64r   c  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   atomic_dec2_32r   h  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   atomic_dec2_64r   m  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   atomic_dec3r   r  r   r   c                 L    t          || d|t          j        j        d           d S r   )r>   r   rC   r   r   s      r   atomic_dec_globalr   w  rr   r   c                 T    t          | |t          j        j        t          d           d S r   )r<   r   rC   r   r   r   s     r   atomic_dec_global_2r   {  r   r   c           	      b    t          | ||t          dt          j        j        t
                     d S r   )r0   r   r   rC   exchr   r   s      r   atomic_exchr     s8    sCfb#{/1AC C C C Cr   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   atomic_exch2r     r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r9   r   r   rC   r   r   r   s     r   atomic_exch3r     r   r   c                 L    t          || d|t          j        j        d           d S r   )r>   r   rC   r   r   s      r   atomic_exch_globalr     r   r   c                     t          d                              |           }i }t          |t          t          t
          d|           |d         |d         |d         |d         fS )Na  
    def atomic(res, ary):
        tx = cuda.threadIdx.x
        bx = cuda.blockIdx.x
        {func}(res, 0, ary[tx, bx])

    def atomic_double_normalizedindex(res, ary):
        tx = cuda.threadIdx.x
        bx = cuda.blockIdx.x
        {func}(res, 0, ary[tx, uint64(bx)])

    def atomic_double_oneindex(res, ary):
        tx = cuda.threadIdx.x
        {func}(res, 0, ary[tx])

    def atomic_double_shared(res, ary):
        tid = cuda.threadIdx.x
        smary = cuda.shared.array(32, float64)
        smary[tid] = ary[tid]
        smres = cuda.shared.array(1, float64)
        if tid == 0:
            smres[0] = res[0]
        cuda.syncthreads()
        {func}(smres, 0, smary[tid])
        cuda.syncthreads()
        if tid == 0:
            res[0] = smres[0]
    )func)r   r   r   rC   atomic_double_normalizedindexatomic_double_oneindexatomic_double_shared)r   formatexecr   r   r   )r   fnslds      r   gen_atomic_extreme_funcsr    sy    
  	 	6 
T		7 8 
Bt6BBBGG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     r4t           j                            | |d          |||                   ||<   d S d S Nr3   )r   gridsizerC   compare_and_swapresoldr"   fill_valgids        r   atomic_compare_and_swapr    sN    
)A,,C
SX~~;//CDD	8SXNNC ~r   c                     t          j        d          }|| j        k     r-t           j                            | ||||                   ||<   d S d S r  )r   r  r	  rC   casr  s        r   atomic_cas_1dimr    sF    
)A,,C
SX~~;??3Xs3x@@C ~r   c                     t          j        d          }|d         | j        d         k     rD|d         | j        d         k     r/t           j                            | ||||                   ||<   d S d S d S )N   r   r3   )r   r  r;   rC   r  r  s        r   atomic_cas_2dimr    sk    
)A,,C
1v	!Q#)A,!6!6;??3Xs3x@@C !6!6r   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/ Z0d0 Z1d1 Z2d2 Z3d3 Z4d4 Z5d5 Z6d6 Z7d7 Z8d8 Z9d9 Z:d: Z;d; Z<d< Z=d= Z>d> Z?d? Z@d@ ZAdA ZBdB ZCdC ZDdD ZEdE ZFdF ZGdG ZHdH ZIdI ZJdJ ZKdK ZLdL ZMdM ZNdN ZOdO ZPdP ZQdQ ZRdR ZSdS ZTdT ZUdU ZVdV ZWdW ZXdX ZYdY ZZdZ Z[d[ Z\d\ Z]d] Z^d^ Z_d_ Z`d` Zada Zbdb ZcdddZdde Zedf Zfdg Zgdh Zhdi Zidj Zjdk Zkdl Zldm Zmdn Zndo Zodp Zpdq Zqdr Zrds Zsdt Ztdu Zudv Zvdw Zwdx Zxdy Zydz Zzd{ Z{d| 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                            d           d S )Nr   )supersetUpnprandomseed)self	__class__s    r   r  zTestCudaAtomics.setUp  s.    
	qr   c                    t           j                            ddd                              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<   |                     t          j        ||k                         |                     t          j        ||k                         d S 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:   i2B//66ryAA88::xxzz5$(#455jAAs###:tx(9::?KK#U#H---x"),,,ty!! 	 	AaMMMQMMMMsd{++,,,x4/0011111r   c                    t           j                            ddd                              t           j                                      dd          }|                                }|                                } t          j        d          t                    } |d         |            t          j        d          t                    } |d         |           |                     t          j        ||dz   k                         |                     t          j        ||dz   k                         d S 	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
   i2B//66ryAAII!QOO88::xxzz848$788EE##C((( =)< = =>N O O(i(222sdQh//000x4!834455555r   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           | 
                    t          j        ||dz   k                         d S 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    i2B//66ryAAII!QOOxxzz848$788EE##C(((sdQh//00000r   c                     t           j                            ddd                              t           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<   |                     t          j        ||k                         |                     t          j        ||k                         d S 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  sL   i2B//66rzBB88::xxzz  )) <); < <=M N N$e$S)))5"4556KLLuh'''x"),,,ty!! 	! 	!AaMMMS MMMMsd{++,,,x4/0011111r   c                    t           j                            ddd                              t           j                                      dd          }|                                }|                                } t          j        d          t                    } |d         |            t          j        d          t                    } |d         |           |                     t          j        ||dz   k                         |                     t          j        ||dz   k                         d S 	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	   i2B//66rzBBJJ1aPP88::xxzz948$899:LMM##C(((7"6778OPP!y!(+++sdQh//000x4!834455555r   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           | 
                    t          j        ||dz   k                         d S 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    i2B//66rzBBJJ1aPPxxzz948$899:LMM##C(((sdQh//00000r   Tc                    t           j        rd S t          t          |                                                                                    }t          dd          rIt          j        	                                dk    rd}nd}|r| d}| 
                    | d|           d S |r| 
                    d|           d S | 
                    d	|           d S )
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      	F 4**,,33556677A 	3|''))G33  ('''MMT+++S11111 33S99999nc22222r   c                    t           j                            dddt           j                  }t          j        d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<   t           j                            ||           t           j                            ||           |                     |           |                     |           d S 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  sK   i2Bbh??hr2:&&88::8$(7889JKKsC   8$(7889OPPsH%%%x"),,,sx 	  	 AQLLLCLLLL

T***

$///((111((11111r   c                    t           j                            ddd                              t           j                                      dd          }|                                }|                                } t          j        d          t                    } |d         |            t          j        d          t                    } |d         |           t           j                            ||dz              t           j                            ||dz              |                     |           |                     |           d S 	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   i2B//66rzBBJJ1aPP88::xxzz0$(/001DEE	35tx 4556NOOY)))

TAX...

$(333((111((66666r   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           t           j
                            ||dz              |                     |           d S 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    i2B//66rzBBJJ1aPPxxzz2DH1223FGG		)S!!!

TAX...((33333r   c                    t           j                            dddt           j                  }t          j        dt           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                            ||           t           j                            ||           |                     |d	           |                     |d	           d S )
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  sV   i2Bbh??hr2:&&88::*!DHSMM":;;	&#'DEE	%c"""uc8,,,x"),,,sx 	  	 AQLLLCLLLL

T***

$///((5(AAA(((FFFFFr   c                    t           j                            ddd                              t           j                                      dd          }|                                }|                                }d} t          j        |          t                    } t          j        |          t                    } |d         |            |d         |           t           j                            ||dz              t           j                            ||dz              |                     |d	
           |                     |d	
           d S 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'   i2B//66rzBBJJ1aPP88::xxzz"!DHSMM"<==	&#'FGG	)S!!!!y!(+++

TAX...

$(333((5(AAA(((FFFFFr   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           t           j
                            ||dz              |                     |d	
           d S 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    i2B//66rzBBJJ1aPPxxzz2DH1223MNN		)S!!!

TAX...((5(AAAAAr   c                    t           j                            ddd                              t           j                  }|                                } t          j        d          t                    } |d         |           t          j	        dt           j                  }t          |j                  D ]}|||         xx         dz  cc<   |                     t          j        ||k                         d S 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    i2B//66ryAAxxzz5$(#455jAAs###x"),,,ty!! 	 	AaMMMQMMMMsd{++,,,,,r   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           | 
                    t          j        ||dz
  k                         d S 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      i2B//66ryAAII!QOOxxzz848$788EE##C(((sdQh//00000r   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           | 
                    t          j        ||dz
  k                         d S 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                            ddd                              t           j                  }|                                                    t           j                  } t          j        d          t                    } |d         |           t          j
        dt           j                  }t          |j                  D ]}|||         xx         dz  cc<   |                     t          j        ||k                         d S 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    i2B//66rzBBxxzz  )) <); < <=M N N$e$S)))x"*---ty!! 	! 	!AaMMMS MMMMsd{++,,,,,r   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           | 
                    t          j        ||dz
  k                         d S 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      i2B//66rzBBJJ1aPPxxzz948$899:LMM##C(((sdQh//00000r   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           | 
                    t          j        ||dz
  k                         d S 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                            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                            ||           d S 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    i2Bbh??hr2:&&:DH9::;LMM		%c"""x"*---sx 	  	 AQLLLCLLLL

T*****r   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           t           j
                            ||dz
             d S 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      i2B//66rzBBJJ1aPPxxzz2DH1223FGG		)S!!!

TAX.....r   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           t           j
                            ||dz
             d S 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                            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                            ||           d S 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    i2Bbh??hr2:&&*!DHSMM":;;		%c"""x"*---sx 	  	 AQLLLCLLLL

T*****r   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           t           j
                            ||dz
             d S 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      i2B//66rzBBJJ1aPPxxzz2DH1223MNN		)S!!!

TAX.....r   c                    t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         |           t           j
                            ||dz
             d S 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                            d          }t           j                            ddd                              t           j                  }|                                } t          j        d          t                    } |d         ||           |                                }t          |j
                  D ]}|||         xx         |z  cc<   |                     t          j        ||k                         d S )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    Y&&s++
i2B//66ryAAxxzz7DH677
CC		%j)))xxzzty!! 	( 	(AaMMMZ'MMMMsd{++,,,,,r   c                    t           j                            d          }t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         ||           | 
                    t          j        |||z  k                         d S 	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      Y&&s++
i2B//66ryAAII!QOOxxzz@48$?@@MM##C444sdZ&778899999r   c                    t           j                            d          }t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         ||           | 
                    t          j        |||z  k                         d S 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                            d          }t           j                            dddt           j                  }t           j                            dddt           j                  }d} t	          j        |          t                    } |d         |||           |                                }t          |j	                  D ]}|||         xx         |z  cc<   t           j
                            ||           d S 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    Y&&s++
i2Bbh??i2Bbh??/!DHSMM"344		%c:...xxzzsx 	' 	'AQLLLJ&LLLL

T*****r   c                    t           j                            d          }t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         ||           t           j
                            |||z             d S 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      Y&&s++
i2B//66ryAAII!QOOxxzz9DH899:MNN		)S*---

TJ%677777r   c                 B   t           j                            d          }t           j                            ddd                              t           j                  }|                                } t          j        d          t                    } |d         ||           t          j	        dt           j                  }t          |j                  D ]}|||         xx         |z  cc<   |                     t          j        ||k                         d S 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    Y&&s++
i2B//66ryAAxxzz7DH677	BB		%j)))x"),,,ty!! 	( 	(AaMMMZ'MMMMsd{++,,,,,r   c                    t           j                            d          }t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         ||           | 
                    t          j        |||z  k                         d S r  )r  r  r'  r(  r   r9  r)  r   r*  r   r-  r.  r  s        r   test_atomic_or2zTestCudaAtomics.test_atomic_or2H      Y&&s++
i2B//66ryAAII!QOOxxzz@48$?@@LL##C444sdZ&778899999r   c                    t           j                            d          }t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         ||           | 
                    t          j        |||z  k                         d S 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                            d          }t           j                            dddt           j                  }t           j                            dddt           j                  }d} t	          j        |          t                    } |d         |||           |                                }t          |j	                  D ]}|||         xx         |z  cc<   t           j
                            ||           d S 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    Y&&s++
i2Bbh??i2Bbh??/!DHSMM"233		%c:...xxzzsx 	' 	'AQLLLJ&LLLL

T*****r   c                    t           j                            d          }t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         ||           t           j
                            |||z             d S 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    Y&&s++
i2B//66ryAAII!QOOxxzz9DH899:LMM		)S*---

TJ%677777r   c                 B   t           j                            d          }t           j                            ddd                              t           j                  }|                                } t          j        d          t                    } |d         ||           t          j	        dt           j                  }t          |j                  D ]}|||         xx         |z  cc<   |                     t          j        ||k                         d S 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    Y&&s++
i2B//66ryAAxxzz7DH677
CC		%j)))x"),,,ty!! 	( 	(AaMMMZ'MMMMsd{++,,,,,r   c                    t           j                            d          }t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         ||           | 
                    t          j        |||z  k                         d S 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                            d          }t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         ||           | 
                    t          j        |||z  k                         d S 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                            d          }t           j                            dddt           j                  }t           j                            dddt           j                  }|                                }d} t          j        |          t                    } |d         |||           t          |j	                  D ]}|||         xx         |z  cc<   t           j
                            ||           d S 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    Y&&s++
i2Bbh??i2Bbh??xxzz/!DHSMM"344		%c:...sx 	' 	'AQLLLJ&LLLL

T*****r   c                    t           j                            d          }t           j                            ddd                              t           j                                      dd          }|                                } t          j        d          t                    } |d         ||           t           j
                            |||z             d S 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                            d|          }t           j                            ddd                              |          }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  sg    ""2e"44y  BR 0077>>)Be,,,tW$$r   c                     t           j                            d|          }t           j                            ddd                              |                              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  s`    ""2U"33y  BR 0077>>FFq!LLt|r   c           	         |                                 } t          j        |          |          }	 |	||f         |||           t          j                            |t          j        ||k    d|dz                        d S 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  x    xxzz!DHSMM$''	#	'7"#Cf555

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

RXdfna%J%JKKKKKr   c           	          |                      t          j                  \  }}}d}|                     ||||ddt                     d S 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  L    #66RY6GG
C2S#z32|LLLLLr   c           	          |                      t          j                  \  }}}d}|                     ||||ddt                     d S 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                     |                      t          j                  \  }}d}|                     |||ddt                     d S Nr  r3   rL   )r  r  r   r  r   r  r  r"   rs  s       r   test_atomic_inc2_32z#TestCudaAtomics.test_atomic_inc2_32  B    11")<<
C)sJQ~FFFFFr   c                     |                      t          j                  \  }}d}|                     |||ddt                     d S 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                     |                      t          j                  \  }}d}|                     |||ddt                     d S r  )r  r  r   r  r   r  s       r   test_atomic_inc3z TestCudaAtomics.test_atomic_inc3  B    11")<<
C)sJQ{CCCCCr   c           	          |                      t          j                  \  }}}d}|                     ||||ddt                     d S r  )r  r  r   r  r   r  s        r   test_atomic_inc_global_32z)TestCudaAtomics.test_atomic_inc_global_32  W    #66RY6GG
C2c3
CB/	1 	1 	1 	1 	1r   c           	          |                      t          j                  \  }}}d}|                     ||||ddt                     d S r  )r  r  r   r  r   r  s        r   test_atomic_inc_global_64z)TestCudaAtomics.test_atomic_inc_global_64  r  r   c                     |                      t          j                  \  }}d}|                     |||ddt                     d S r  )r  r  r   r  r   r  s       r   test_atomic_inc_global_2_32z+TestCudaAtomics.test_atomic_inc_global_2_32  C    11")<<
C)sJQ7JKKKKKr   c                     |                      t          j                  \  }}d}|                     |||ddt                     d S 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                 8   |                                 } t          j        |          |          }	 |	||f         |||           t          j                            |t          j        |dk    |t          j        ||k    ||dz
                                 d S r2   r  r  s
             r   check_dec_indexzTestCudaAtomics.check_dec_index      xxzz!DHSMM$''	#	'7"#Cf555

RXdai.0htf}7=7;ax/A /A&B &B 	C 	C 	C 	C 	Cr   c                 8   |                                 } t          j        |          |          }	 |	||f         |||           t          j                            |t          j        |dk    |t          j        ||k    ||dz
                                 d S r2   r  r  s
             r   check_dec_index2z TestCudaAtomics.check_dec_index2  r  r   c                 6   |                                 } t          j        |          |          } |||f         ||           t          j                            |t          j        |dk    |t          j        ||k    ||dz
                                 d S r2   r  r  s	            r   	check_deczTestCudaAtomics.check_dec  s    xxzz!DHSMM$''	#	'7"#C000

RXdai.0htf}7=7;ax/A /A&B &B 	C 	C 	C 	C 	Cr   c           	          |                      t          j                  \  }}}d}|                     ||||ddt                     d S r  )r  r  r   r  r   r  s        r   test_atomic_dec_32z"TestCudaAtomics.test_atomic_dec_32  r  r   c           	          |                      t          j                  \  }}}d}|                     ||||ddt                     d S r  )r  r  r   r  r   r  s        r   test_atomic_dec_64z"TestCudaAtomics.test_atomic_dec_64  r  r   c                     |                      t          j                  \  }}d}|                     |||ddt                     d S r  )r  r  r   r	  r   r  s       r   test_atomic_dec2_32z#TestCudaAtomics.test_atomic_dec2_32  r  r   c                     |                      t          j                  \  }}d}|                     |||ddt                     d S r  )r  r  r   r	  r   r  s       r   test_atomic_dec2_64z#TestCudaAtomics.test_atomic_dec2_64  r  r   c                     |                      t          j                  \  }}d}|                     |||ddt                     d S r  )r  r  r   r	  r   r  s       r   test_atomic_dec3_newz$TestCudaAtomics.test_atomic_dec3_new  r  r   c           	          |                      t          j                  \  }}}d}|                     ||||ddt                     d S r  )r  r  r   r  r   r  s        r   test_atomic_dec_global_32z)TestCudaAtomics.test_atomic_dec_global_32!  r  r   c           	          |                      t          j                  \  }}}d}|                     ||||ddt                     d S r  )r  r  r   r  r   r  s        r   test_atomic_dec_global_64z)TestCudaAtomics.test_atomic_dec_global_64'  r  r   c                     |                      t          j                  \  }}d}|                     |||ddt                     d S r  )r  r  r   r	  r   r  s       r   test_atomic_dec_global2_32z*TestCudaAtomics.test_atomic_dec_global2_32-  r   r   c                     |                      t          j                  \  }}d}|                     |||ddt                     d S 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                            ddt           j                  }t           j                            ddd                              t           j                  }t          j        dt           j                  } t          j        d          t                    } |d         |||           t           j	        
                    ||           d S )	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    Y&&r3bi&@@
i2B//66ryAAi"),,,BDHABB;OO		%c:...

Z00000r   c                    t           j                            ddt           j                  }t           j                            ddd                              t           j                                      dd          } t          j        d	          t                    } |d
         ||           t           j	        
                    ||           d S )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      Y&&r3bi&@@
i2B//66ryAAII!QOO9DH899,GG		)S*---

Z00000r   c                    t           j                            ddt           j                  }t           j                            ddd                              t           j                                      dd          } t          j        d	          t                    } |d
         ||           t           j	        
                    ||           d S )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                            ddt           j                  }t          j        dt           j                  }t           j                            dddt           j                  }d} t          j        |          t                    } |d         |||           t           j        	                    ||           d S )	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    Y&&r3bi&@@
i"),,,i2Bbi@@2!DHSMM"455		%c:...

Z00000r   c                 d   t           j                            ||d                              |          }t          j        d|j                  }t          j        t                    } |d         ||           t          j	        |          }t           j
                            ||           d S )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    y  Rh 77>>uEEhq
+++HZ((		&#t$$$vd||

T*****r   c                 J    |                      t          j        dd           d S N   r&  r-  r.  )r0  r  r  r  s    r   test_atomic_max_int32z%TestCudaAtomics.test_atomic_max_int32c  %    BHEBBBBBr   c                 J    |                      t          j        dd           d S Nr   r4  r5  )r0  r  r   r6  s    r   test_atomic_max_uint32z&TestCudaAtomics.test_atomic_max_uint32f  %    BI!>>>>>r   c                 J    |                      t          j        dd           d S r2  )r0  r  rb  r6  s    r   test_atomic_max_int64z%TestCudaAtomics.test_atomic_max_int64i  r8  r   c                 J    |                      t          j        dd           d S r:  )r0  r  r   r6  s    r   test_atomic_max_uint64z&TestCudaAtomics.test_atomic_max_uint64l  r<  r   c                 J    |                      t          j        dd           d S r2  )r0  r  r   r6  s    r   test_atomic_max_float32z'TestCudaAtomics.test_atomic_max_float32o  %    BJ6eDDDDDr   c                 J    |                      t          j        dd           d S r2  )r0  r  r   r6  s    r   test_atomic_max_doublez&TestCudaAtomics.test_atomic_max_doubler  rC  r   c                    t           j                            ddd                              t           j                  }t          j        dt           j                  } t          j        d          t                    } |d         ||           t          j	        |          }t           j
                            ||           d S 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    y  E 99@@LLhq"*%%>DH=>>-/ /		&#t$$$vd||

T*****r   c                    t           j                            ddd                              t           j                  }t          j        dt           j                  } t          j        d          t                    } |d         ||           t          j	        |          }t           j
                            ||           d S 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    y  Cb 1188DDhq"*%%<DH;<<&( (		%d###vd||

T*****r   c                 f   t           j                            ||d                              |          }t          j        dg|j                  }t          j        t                    } |d         ||           t          j	        |          }t           j
                            ||           d S )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    y  Rh 77>>uEEhwdj111HZ((		&#t$$$vd||

T*****r   c                 J    |                      t          j        dd           d S r2  )rU  r  r  r6  s    r   test_atomic_min_int32z%TestCudaAtomics.test_atomic_min_int32  r8  r   c                 J    |                      t          j        dd           d S r:  )rU  r  r   r6  s    r   test_atomic_min_uint32z&TestCudaAtomics.test_atomic_min_uint32  r<  r   c                 J    |                      t          j        dd           d S r2  )rU  r  rb  r6  s    r   test_atomic_min_int64z%TestCudaAtomics.test_atomic_min_int64  r8  r   c                 J    |                      t          j        dd           d S r:  )rU  r  r   r6  s    r   test_atomic_min_uint64z&TestCudaAtomics.test_atomic_min_uint64  r<  r   c                 J    |                      t          j        dd           d S r2  )rU  r  r   r6  s    r   test_atomic_min_floatz%TestCudaAtomics.test_atomic_min_float  rC  r   c                 J    |                      t          j        dd           d S r2  )rU  r  r   r6  s    r   test_atomic_min_doublez&TestCudaAtomics.test_atomic_min_double  rC  r   c                    t           j                            ddd                              t           j                  }t          j        dt           j                  dz  } t          j        d          t                    } |d         ||           t          j	        |          }t           j
                            ||           d S 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    y  E 99@@LLga$$u,>DH=>>-/ /		&#t$$$vd||

T*****r   c                    t           j                            ddd                              t           j                  }t          j        dt           j                  dz  } t          j        d          t                    } |d         ||           t          j	        |          }t           j
                            ||           d S 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    y  Cb 1188DDga$$s*<DH;<<&( (		%d###vd||

T*****r   c                     t          j        d          |          }t          j                            ddd                              t          j                  }t          j        dt          j                  t          j        z   } |d         ||           t          j	        
                    |t          j        g           d S )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    >DH=>>tDD	y  Ce 44;;BJGGhq"*%%.	$T"""

bfX.....r   c                     t          j        d          |          }t          j                            ddd                              t          j                  }|                                }t          j        dt          j                  t          j	        z   } |d         ||           t          j
                            ||           d S )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    >DH=>>tDD	i3Q//66rzBBxxzzx
++bf4	$T"""

T*****r   c                 :    |                      t                     d S r   )rl  rS  r6  s    r   test_atomic_min_nan_locationz,TestCudaAtomics.test_atomic_min_nan_location      --j99999r   c                 :    |                      t                     d S r   )rl  r*  r6  s    r   test_atomic_max_nan_locationz,TestCudaAtomics.test_atomic_max_nan_location  rq  r   c                 :    |                      t                     d S r   )rn  rS  r6  s    r   test_atomic_min_nan_valz'TestCudaAtomics.test_atomic_min_nan_val      ((44444r   c                 :    |                      t                     d S r   )rn  r*  r6  s    r   test_atomic_max_nan_valz'TestCudaAtomics.test_atomic_max_nan_val  rv  r   c                    t           j                            ddd                              t           j                  }t          j        dt           j                  }d} t          j        |          t                    } |d         ||           t          j	        |          }t           j
                            ||           d S 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    y  BR 0077
CChq"*%%,!DHSMM":;;		%d###vd||

T*****r   c                    t           j                            ddd                              t           j                  }t          j        dt           j                  dz  }d} t          j        |          t                    } |d         ||           t          j	        |          }t           j
                            ||           d S 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    y  BR 0077
CCga$$r),!DHSMM":;;		%d###vd||

T*****r   r3   c                    |g|dz  z  |g|dz  z  z   }t           j                            |           t          j        ||          }|dk    rd|_        t          j        |          }t           j                            dd|j                                      |j                  }	||k    }
||k    }t          j        |          }|	|
         ||
<   |||<   |	                                }t          j        |          }|dk    r |d         |||	|           n |d         |||	|           t           j                            ||           t           j                            ||           d S )	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  sg   fQ6(a1f"55
	#jE***199 CImC  i2CI66==ciHH4K	Vm]3''
 #I
9"(
;XXZZ
HX&&	199Ifc3T2222)I()#sC>>>

%%j#666

%%j#66666r   c                 X    |                      dddt          j        t                     d S 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  4    3r 7 	 	9 	9 	9 	9 	9r   c                 X    |                      dddt          j        t                     d S 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                            ddt           j                  }t           j                            ddt           j                  }|                     d||t           j        t
                     d S 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  p    	!!"c!;;)##Ar#;;5	 7 	 	9 	9 	9 	9 	9r   c                    t           j                            ddt           j                  }t           j                            ddt           j                  }|                     d||t           j        t
                     d S r  )r  r  r'  r   r  r  r  s      r   test_atomic_compare_and_swap4z-TestCudaAtomics.test_atomic_compare_and_swap4  r  r   c                 X    |                      dddt          j        t                     d S r  )r  r  r  r  r6  s    r   test_atomic_cas_1dimz$TestCudaAtomics.test_atomic_cas_1dim%  4    3r / 	 	1 	1 	1 	1 	1r   c                 Z    |                      dddt          j        t          d           d S )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)  6    3r /a 	 	9 	9 	9 	9 	9r   c                 X    |                      dddt          j        t                     d S r  )r  r  rb  r  r6  s    r   test_atomic_cas2_1dimz%TestCudaAtomics.test_atomic_cas2_1dim-  r  r   c                 Z    |                      dddt          j        t          d           d S )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                            ddt           j                  }t           j                            ddt           j                  }|                     d||t           j        t
                     d S r  )r  r  r'  r   r  r  r  s      r   test_atomic_cas3_1dimz%TestCudaAtomics.test_atomic_cas3_1dim5  p    	!!"c!;;)##Ar#;;5	 / 	 	1 	1 	1 	1 	1r   c                 
   t           j                            ddt           j                  }t           j                            ddt           j                  }|                     d||t           j        t
          d           d S 	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;  r    	!!"c!;;)##Ar#;;5	 /a 	 	9 	9 	9 	9 	9r   c                    t           j                            ddt           j                  }t           j                            ddt           j                  }|                     d||t           j        t
                     d S r  )r  r  r'  r   r  r  r  s      r   test_atomic_cas4_1dimz%TestCudaAtomics.test_atomic_cas4_1dimA  r  r   c                 
   t           j                            ddt           j                  }t           j                            ddt           j                  }|                     d||t           j        t
          d           d S r  )r  r  r'  r   r  r  r  s      r   test_atomic_cas4_2dimz%TestCudaAtomics.test_atomic_cas4_2dimG  r  r   c                 0   t          j        dt           j                  }||d<    |d         |           t          j        |          r/|                     t          j        |d                              d S |                     |d         |           d S )Nr  r%  r   rj  r3   )r  r+  r   isnanr-  assertEqualr  rZ  initialr   s       r   _test_atomic_returns_oldz(TestCudaAtomics._test_atomic_returns_oldR  s    HQbj)))!tQ8G 	,OOBHQqTNN+++++QqT7+++++r   c                 \    t           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S r2   )r   rC   rD   r   s    r   rZ  z;TestCudaAtomics.test_atomic_add_returns_old.<locals>.kernel\  !    ;??1a++AaDDDr   r  r   r*  r  r  rZ  s     r   test_atomic_add_returns_oldz+TestCudaAtomics.test_atomic_add_returns_old[  ;    		, 	, 
	, 	%%fb11111r   c                 \    t           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S 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           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S Nr   r  r3   r  r  s    r   rZ  zCTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernelj  !    ;??1a,,AaDDDr   r3   r  r  s     r   #test_atomic_max_returns_old_replacez3TestCudaAtomics.test_atomic_max_returns_old_replacei  s;    		- 	- 
	- 	%%fa00000r   c                 p    t           j        d             }|                     |t          j                   d S )Nc                 N    t           j                            | dd          | d<   d S 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=    		, 	, 
	, 	%%fbf55555r   c                 \    t           j        d             }|                     |d           d S )Nc                 b    t           j                            | dt          j                  | d<   d S r2   )r   rC   r+  r  rk  r  s    r   rZ  zCTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernelx  #    ;??1a00AaDDDr   r  r  r  s     r   #test_atomic_max_returns_old_nan_valz3TestCudaAtomics.test_atomic_max_returns_old_nan_valw  ;    		1 	1 
	1 	%%fb11111r   c                 \    t           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S 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11111r   c                 \    t           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S 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             }|                     |t          j                   d S )Nc                 N    t           j                            | dd          | d<   d S 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=    		- 	- 
	- 	%%fbf55555r   c                 \    t           j        d             }|                     |d           d S )Nc                 b    t           j                            | dt          j                  | d<   d S 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                 t   t           j                            ||d                              |          }||dd d<   t          j        d|j                  }t          j        t                    } |d         ||           t          j	        |          }t           j
                            ||           d S )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    y  Rh 77>>uEEQTT
hq
+++H]++		&#t$$$y

T*****r   c                 L    |                      t          j        ddd           d S Nr3  r4  r   r&  r-  r.  r  )r  r  r  r6  s    r   test_atomic_nanmax_int32z(TestCudaAtomics.test_atomic_nanmax_int32  4      rxFu*+ 	! 	- 	- 	- 	- 	-r   c                 L    |                      t          j        ddd           d S Nr   r4  r  )r  r  r   r6  s    r   test_atomic_nanmax_uint32z)TestCudaAtomics.test_atomic_nanmax_uint32  4      ryQ5*+ 	! 	- 	- 	- 	- 	-r   c                 L    |                      t          j        ddd           d S r  )r  r  rb  r6  s    r   test_atomic_nanmax_int64z(TestCudaAtomics.test_atomic_nanmax_int64  r  r   c                 L    |                      t          j        ddd           d S r  )r  r  r   r6  s    r   test_atomic_nanmax_uint64z)TestCudaAtomics.test_atomic_nanmax_uint64  r  r   c                 `    |                      t          j        ddt          j                   d S Nr3  r4  r  )r  r  r   rk  r6  s    r   test_atomic_nanmax_float32z*TestCudaAtomics.test_atomic_nanmax_float32  6      rzf*,& 	! 	2 	2 	2 	2 	2r   c                 `    |                      t          j        ddt          j                   d S r  )r  r  r   rk  r6  s    r   test_atomic_nanmax_doublez)TestCudaAtomics.test_atomic_nanmax_double  r  r   c                    t           j                            ddd                              t           j                  }t           j        |dd d<   t          j        dg|j                  }d} t          j	        |          t                    } |d         ||           t          j        |          }t           j                            ||           d S 	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    y  BR 0077
CCVQTT
hs$*---,!DHSMM"=>>		%d###y

T*****r   c                    t           j                            ddd                              t           j                  }t           j        |dd d<   t          j        dt           j                  } t          j        d          t                    } |d         ||           t          j
        |          }t           j                            ||           d S 	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    y  Cb 1188DDVQTT
hq"*%%<DH;<<&( (		%d###y

T*****r   c                 v   t           j                            ||d                              |          }||dd d<   t          j        dg|j                  }t          j        t                    } |d         ||           t          j	        |          }t           j
                            ||           d S )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    y  Rh 77>>uEEQTT
hwdj111H]++		&#t$$$y

T*****r   c                 L    |                      t          j        ddd           d S r  )r  r  r  r6  s    r   test_atomic_nanmin_int32z(TestCudaAtomics.test_atomic_nanmin_int32  r  r   c                 L    |                      t          j        ddd           d S r  )r  r  r   r6  s    r   test_atomic_nanmin_uint32z)TestCudaAtomics.test_atomic_nanmin_uint32  r  r   c                 L    |                      t          j        ddd           d S r  )r  r  rb  r6  s    r   test_atomic_nanmin_int64z(TestCudaAtomics.test_atomic_nanmin_int64  r  r   c                 L    |                      t          j        ddd           d S r  )r  r  r   r6  s    r   test_atomic_nanmin_uint64z)TestCudaAtomics.test_atomic_nanmin_uint64  r  r   c                 `    |                      t          j        ddt          j                   d S r  )r  r  r   rk  r6  s    r   test_atomic_nanmin_floatz(TestCudaAtomics.test_atomic_nanmin_float  r  r   c                 `    |                      t          j        ddt          j                   d S r  )r  r  r   rk  r6  s    r   test_atomic_nanmin_doublez)TestCudaAtomics.test_atomic_nanmin_double  r  r   c                    t           j                            ddd                              t           j                  }t           j        |dd d<   t          j        dg|j                  }d} t          j	        |          t                    } |d         ||           t          j        |          }t           j                            ||           d S 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    y  BR 0077
CCVQTT
ht4:...,!DHSMM"=>>		%d###y

T*****r   c                    t           j                            ddd                              t           j                  }t           j        |dd d<   t          j        dgt           j                  } t          j        d          t                    } |d         ||           t          j
        |          }t           j                            ||           d S 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    y  Cb 1188DDVQTT
hubj))<DH;<<&( (		%d###y

T*****r   c                    t          j        dt           j                  }||d<   t           j        |d<    |d         |           t          j        |          r\|                     t          j        |d                              |                     t          j        |d                              d S |                     |d         |           d S )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    HQbj)))!v!tQ8G 	,RXad^^,,,OOBHQqTNN+++++QqT7+++++r   c                 \    t           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S r2   r   rC   r  r  s    r   rZ  zITestCudaAtomics.test_atomic_nanmax_returns_old_no_replace.<locals>.kernel  #    ;%%aA..AaDDDr   r  r   r*  r%  r  s     r   )test_atomic_nanmax_returns_old_no_replacez9TestCudaAtomics.test_atomic_nanmax_returns_old_no_replace  s;    		/ 	/ 
	/ 	))&"55555r   c                 \    t           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S r  r(  r  s    r   rZ  zFTestCudaAtomics.test_atomic_nanmax_returns_old_replace.<locals>.kernel"  #    ;%%aB//AaDDDr   r3   r*  r  s     r   &test_atomic_nanmax_returns_old_replacez6TestCudaAtomics.test_atomic_nanmax_returns_old_replace!  s;    		0 	0 
	0 	))&!44444r   c                 p    t           j        d             }|                     |t          j                   d S )Nc                 N    t           j                            | dd          | d<   d S 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=    		/ 	/ 
	/ 	))&"&99999r   c                 \    t           j        d             }|                     |d           d S )Nc                 b    t           j                            | dt          j                  | d<   d S r2   )r   rC   r  r  rk  r  s    r   rZ  zFTestCudaAtomics.test_atomic_nanmax_returns_old_nan_val.<locals>.kernel0  %    ;%%aBF33AaDDDr   r  r*  r  s     r   &test_atomic_nanmax_returns_old_nan_valz6TestCudaAtomics.test_atomic_nanmax_returns_old_nan_val/  ;    		4 	4 
	4 	))&"55555r   c                 \    t           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S 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 
	0 	))&"55555r   c                 \    t           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S 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             }|                     |t          j                   d S )Nc                 N    t           j                            | dd          | d<   d S 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 
	0 	))&"&99999r   c                 \    t           j        d             }|                     |d           d S )Nc                 b    t           j                            | dt          j                  | d<   d S 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 2 2$6 6 61 1 12 2 2$6 6 61 1 13 3 3 322 2 2(7 7 7 4 4 4G G G*G G G"B B B
- 
- 
-1 1 11 1 1
- 
- 
-1 1 11 1 1
+ 
+ 
+/ / // / /+ + +/ / // / /- - -: : :: : :+ + +8 8 8- - -: : :: : :+ + +8 8 8- - -: : :: : :+ + +8 8 8% % %  
L L LL L LL L LM M M
M M M
G G G
G G G
D D D
1 1 11 1 1L L L
L L L
C C CC C CC C CM M M
M M M
G G G
G G G
D D D
1 1 11 1 1L L L
L L L
1 1 11 1 11 1 11 1 1+ + +C C C? ? ?C C C? ? ?E E EE E E+ + ++ + ++ + +C C C? ? ?C C C? ? ?E E EE E E+ + ++ + +*/ / /+ + +: : :: : :5 5 55 5 5+ + ++ + +7 7 7 769 9 99 9 99 9 99 9 91 1 19 9 91 1 19 9 91 1 19 9 91 1 19 9 9, , ,2 2 22 2 21 1 16 6 62 2 22 2 22 2 26 6 62 2 2+ + +- - -- - -- - -- - -2 2 22 2 2	+ 	+ 	+	+ 	+ 	++ + +- - -- - -- - -- - -2 2 22 2 2	+ 	+ 	+	+ 	+ 	+	, 	, 	,6 6 65 5 5: : :6 6 66 6 66 6 6: : :6 6 6 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     sv             8 8 8 8 8 8 8 8 8 8 8 8 8 8 D D D D D D D D D D       
   
   
   
   
	 	 	 
   
   
  K K K
J J J
H H H
G G G
M M M
O O O
N N N
H H H
G G G
M M M
I I IH H HP P PO O O% % %
M M M
L L L
H H H
G G G
M M M
K K K
H H H
M M M
O O O
H H H
M M M
M M M
H H H
M M M
I I I% % %
M M M
L L L
I I I
N N N
J J J7 7 7
K K K
H H H
M M M
I I I7 7 7
K K K
H H H
M M M
I I I7 7 7
B B B
D D D
H H H
H H H
M M M
I I I7 7 7
B B B
D D D
H H H
H H H
M M M
I I I7 7 7
C C C
I I I
I I I
J J J!F !F !FJ 656GHH.0J556GHH.0J 122=4 ; 122=4 ;O O OA A AA A A{6 {6 {6 {6 {6l {6 {6 {6|# zHMOOOOO r   