
    ܙd                     b   d dl Zd dlmZ d dlmZmZmZmZm	Z	 d dl
mZmZmZmZ d dlmZ  ej        d          d             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& Z0d' Z1d( Z2d) Z3d* Z4d+ Z5d, Z6d- Z7d. Z8d/ Z9d0 Z:d1 Z;d2 Z<d3 Z=d4 Z>d5 Z?d6 Z@d7 ZAd8 ZBd9 ZCd: ZDd; ZEd< ZFd= ZGd> ZHd? ZId@ ZJdA ZKdB ZLdC ZMdD ZNdE ZOdF ZPdG ZQdH ZRdI ZSdJ ZTdK ZUdL ZVdM ZWdN ZXdO ZYdP ZZdQ Z[ e[dR          \  Z\Z]Z^Z_ e[dS          \  Z`ZaZbZc e[dT          \  ZdZeZfZg e[dU          \  ZhZiZjZkdV ZldW ZmdX Zn G dY dZe          Zoepd[k    r ejq                     dS dS )\    N)dedent)cudauint32uint64float32float64)unittestCUDATestCaseskip_unless_cc_50cc_X_or_above)configT)devicec                      t          |           S N)r   nums    Dlib/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_shared2r1   )   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_sharedr:   6   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 r3   )r   r   r   r5   shape)r#   r%   r(   r7   r+   r8   r9   r.   s           r   atomic_binary_2dim_globalr=   F   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?   P   sP     .
C
c#h&
'
'C "M!JsCr   c                 f    t          | | dt          dt          j        j        t
          dd	  	         d S Nr4       r   Fr/   r   r   atomicaddr   r#   s    r   
atomic_addrG   Z   ;    c362"ko/?EK K K K Kr   c                 f    t          | | dt          dt          j        j        t
          dd	  	         d S )Nr4   rB   r   TrC   rF   s    r   atomic_add_wraprJ   _   s;    c362"ko/?DJ J J J Jr   c           	      b    t          | dt          dt          j        j        t
          d           d S Nr4         Fr:   r   r   rD   rE   r   rF   s    r   atomic_add2rQ   d   7    c1ff"ko/?H H H H Hr   c           	      b    t          | dt          dt          j        j        t
          d           d S )Nr4   rM   TrP   rF   s    r   atomic_add2_wraprT   i   s7    c1ff"ko/?G G G G Gr   c           	      b    t          | dt          dt          j        j        t
          d           d S rL   )r:   r   r   rD   rE   r   rF   s    r   atomic_add3rV   n   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      ?rB           Fr/   r   r   rD   rE   r   rF   s    r   atomic_add_floatr]   s   ;    c3Wb"ko/A3O O O O Or   c                 f    t          | | dt          dt          j        j        t
          dd	  	         d S NrZ   rB   r[   Tr\   rF   s    r   atomic_add_float_wrapra   x   s;    c3Wb"ko/A3N N N N Nr   c           	      b    t          | dt          dt          j        j        t
          d           d S NrZ   rM   Fr:   r   r   rD   rE   r   rF   s    r   atomic_add_float_2re   }   7    c3"ko/?H H H H Hr   c           	      b    t          | dt          dt          j        j        t
          d           d S NrZ   rM   Trd   rF   s    r   atomic_add_float_2_wrapri      7    c3"ko/?G G G G Gr   c           	      b    t          | dt          dt          j        j        t
          d           d S rc   )r:   r   r   rD   rE   r   rF   s    r   atomic_add_float_3rl      7    c3"ko/DeM M M M Mr   c                 L    t          || ddt          j        j        d           d S NrB   rZ   Fr?   r   rD   rE   r$   r#   s     r   atomic_add_double_globalrr      $    c3C%HHHHHr   c                 L    t          || ddt          j        j        d           d S )NrB   rZ   Trp   rq   s     r   atomic_add_double_global_wrapru      s$    c3C$GGGGGr   c                 T    t          | dt          j        j        t          d           d S Nr4   Fr=   r   rD   rE   r   rF   s    r   atomic_add_double_global_2ry      s#    c1dko7GOOOOOr   c                 T    t          | dt          j        j        t          d           d S )Nr4   Trx   rF   s    r   atomic_add_double_global_2_wrapr{      s#    c1dko7GNNNNNr   c                 T    t          | dt          j        j        t          d           d S rw   )r=   r   rD   rE   r   rF   s    r   atomic_add_double_global_3r}      s.    c1dko7L#% % % % %r   c                 f    t          || dt          dt          j        j        t
          dd	  	         d S rY   r/   r   r   rD   rE   r   rq   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   rq   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 rc   r:   r   r   rD   rE   r   rF   s    r   atomic_add_double_2r      rf   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rh   r   rF   s    r   atomic_add_double_2_wrapr      rj   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rc   )r:   r   r   rD   rE   r   rF   s    r   atomic_add_double_3r      rm   r   c                 f    t          | | dt          dt          j        j        t
          dd	  	         d S rA   )r/   r   r   rD   subr   rF   s    r   
atomic_subr      rH   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rL   )r:   r   r   rD   r   r   rF   s    r   atomic_sub2r      rR   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rL   )r:   r   r   rD   r   r   rF   s    r   atomic_sub3r      rW   r   c                 f    t          | | dt          dt          j        j        t
          dd	  	         d S rY   )r/   r   r   rD   r   r   rF   s    r   atomic_sub_floatr      r^   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rc   )r:   r   r   rD   r   r   rF   s    r   atomic_sub_float_2r      rf   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rc   )r:   r   r   rD   r   r   rF   s    r   atomic_sub_float_3r      rm   r   c                 f    t          || dt          dt          j        j        t
          dd	  	         d S rY   )r/   r   r   rD   r   r   rq   s     r   atomic_sub_doubler      r   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rc   )r:   r   r   rD   r   r   rF   s    r   atomic_sub_double_2r      rf   r   c           	      b    t          | dt          dt          j        j        t
          d           d S rc   r:   r   r   rD   r   r   rF   s    r   atomic_sub_double_3r      rm   r   c                 L    t          || ddt          j        j        d           d S ro   )r?   r   rD   r   rq   s     r   atomic_sub_double_globalr      rs   r   c                 T    t          | dt          j        j        t          d           d S )NrZ   F)r=   r   rD   r   r   rF   s    r   atomic_sub_double_global_2r      s.    c39I#% % % % %r   c           	      b    t          | dt          dt          j        j        t
          d           d S rc   r   rF   s    r   atomic_sub_double_global_3r      rm   r   c                 f    t          | | |t          dt          j        j        t
          dd	  	         d S )NrB   r4   F)r/   r   r   rD   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 NrM   F)r:   r   r   rD   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   )r:   r   r   rD   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 NrB   F)r?   r   rD   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   rD   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 NrB   r   F)r/   r   r   rD   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   )r:   r   r   rD   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   )r:   r   r   rD   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   rD   r   r   s      r   atomic_or_globalr     rs   r   c                 T    t          | |t          j        j        t          d           d S r   )r=   r   rD   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   rD   xorr   r   s     r   
atomic_xorr   %  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r:   r   r   rD   r   r   r   s     r   atomic_xor2r   *  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r:   r   r   rD   r   r   r   s     r   atomic_xor3r   /  r   r   c                 L    t          || d|t          j        j        d           d S r   )r?   r   rD   r   r   s      r   atomic_xor_globalr   4  rs   r   c                 T    t          | |t          j        j        t          d           d S r   )r=   r   rD   r   r   r   s     r   atomic_xor_global_2r   8  r   r   c           	      b    t          | ||t          dt          j        j        t
                     d S NrB   )r1   r   r   rD   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   )r1   r   r   rD   r   r   r   s      r   atomic_inc64r   B  7    sCfb#{0BD D D D Dr   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r:   r   r   rD   r   r   r   s     r   atomic_inc2_32r   G  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r:   r   r   rD   r   r   r   s     r   atomic_inc2_64r   L  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r:   r   r   rD   r   r   r   s     r   atomic_inc3r   Q  r   r   c                 L    t          || d|t          j        j        d           d S r   )r?   r   rD   r   r   s      r   atomic_inc_globalr   V  rs   r   c                 T    t          | |t          j        j        t          d           d S r   )r=   r   rD   r   r   r   s     r   atomic_inc_global_2r   Z  r   r   c           	      b    t          | ||t          dt          j        j        t
                     d S r   )r1   r   r   rD   decr   r   s      r   atomic_dec32r   _  r   r   c           	      b    t          | ||t          dt          j        j        t
                     d S r   )r1   r   r   rD   r   r   r   s      r   atomic_dec64r   d  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r:   r   r   rD   r   r   r   s     r   atomic_dec2_32r   i  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r:   r   r   rD   r   r   r   s     r   atomic_dec2_64r   n  r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r:   r   r   rD   r   r   r   s     r   atomic_dec3r   s  r   r   c                 L    t          || d|t          j        j        d           d S r   )r?   r   rD   r   r   s      r   atomic_dec_globalr   x  rs   r   c                 T    t          | |t          j        j        t          d           d S r   )r=   r   rD   r   r   r   s     r   atomic_dec_global_2r   |  r   r   c           	      b    t          | ||t          dt          j        j        t
                     d S r   )r1   r   r   rD   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   )r:   r   r   rD   r   r   r   s     r   atomic_exch2r     r   r   c           	      b    t          | |t          dt          j        j        t
          d           d S r   )r:   r   r   rD   r   r   r   s     r   atomic_exch3r     r   r   c                 L    t          || d|t          j        j        d           d S r   )r?   r   rD   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   rD   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 Nr4   )r   gridsizerD   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
  rD   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   r4   )r   r	  r<   rD   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
ed
             Zd Zd Zed             Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Z d Z!d  Z"d! Z#d" Z$d# Z%d$ Z&d% Z'd& Z(d' Z)d( Z*d) Z+d* Z,d+ Z-d, Z.d- Z/d. Z0d/ Z1d0 Z2d1 Z3d2 Z4d3 Z5d4 Z6d5 Z7d6 Z8d7 Z9d8 Z:d9 Z;d: Z<d; Z=d< Z>d= Z?d> Z@d? ZAd@ ZBdA ZCdB ZDdC ZEdD ZFdE ZGdF ZHdG ZIdH ZJdI ZKdJ ZLdK ZMdL ZNdM ZOdN ZPdO ZQdP ZRdQ ZSdR ZTdS ZUdT ZVdU ZWdV ZXdW ZYdX ZZdY Z[dZ Z\d[ Z]d\ Z^d] Z_d^ Z`d_ Zad` Zbda Zcdb Zdddd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 Zd 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   rB   r
  zvoid(uint32[:])r4   rB   dtyper4   )r  r  randintastyper   copyr   jitrG   rJ   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   rB   r$  rN   rO   zvoid(uint32[:,:])r4   rM   r4   )r  r  r(  r)  r   reshaper*  r   r+  rQ   rT   r.  r/  )r   r#   r0  r1  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 r8  )r  r  r(  r)  r   r:  r*  r   r+  rV   r.  r/  r   r#   r1  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   rB   r$  zvoid(float32[:])r%  r&  rZ   )r  r  r(  r)  r   r*  intpr   r+  r]   ra   r,  r   r-  r
  r.  r/  )r   r#   r0  r1  cuda_atomic_add_floatadd_float_wrapr4  r5  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   rB   r$  rN   rO   zvoid(float32[:,:])r9  r4   )r  r  r(  r)  r   r:  r*  r   r+  re   ri   r.  r/  )r   r#   r0  r1  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 rI  )r  r  r(  r)  r   r:  r*  r   r+  rl   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          r2|r|                     d|           d S |                     d|           d S |r|                     d|           d S |                     d|           d S )N   r   zatom.shared.add.f64zatom.add.f64zatom.shared.cas.b64zatom.cas.b64)r   ENABLE_CUDASIMnextiterinspect_asmvaluesr   assertIn)r   kernelr    asms       r   assertCorrectFloat64Atomicsz+TestCudaAtomics.assertCorrectFloat64Atomics+  s      	F 4**,,33556677A 		3 33S99999nc22222 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   rB   r
  r'  void(int64[:], float64[:])r%  r&  rZ   )r  r  r(  int64r,  r   r*  r   r+  r   r   r   r-  r
  testingassert_equalrX  )r   r$   r#   r0  cuda_fnwrap_fnr4  r5  s           r   test_atomic_add_doublez&TestCudaAtomics.test_atomic_add_double<  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   rB   r$  rN   rO   void(float64[:,:])r9  r4   )r  r  r(  r)  r   r:  r*  r   r+  r   r   r^  r_  rX  )r   r#   r0  r1  r`  cuda_fn_wraps         r   test_atomic_add_double_2z(TestCudaAtomics.test_atomic_add_double_2Q  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 rd  )r  r  r(  r)  r   r:  r*  r   r+  r   r^  r_  rX  r   r#   r1  	cuda_funcs       r   test_atomic_add_double_3z(TestCudaAtomics.test_atomic_add_double_3a  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   rB   r[  r\  r%  r&  rZ   Fr    )r  r  r(  r]  r,  r   r*  r   r+  rr   ru   r   r-  r
  r^  r_  rX  )	r   r$   r#   r0  sigrj  wrap_cuda_funcr4  r5  s	            r   test_atomic_add_double_globalz-TestCudaAtomics.test_atomic_add_double_globalj  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   rB   r$  rN   rO   re  r9  r4   Frm  )r  r  r(  r)  r   r:  r*  r   r+  ry   r{   r^  r_  rX  )r   r#   r0  r1  rn  rj  ro  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 rr  )r  r  r(  r)  r   r:  r*  r   r+  r}   r^  r_  rX  ri  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#   r1  cuda_atomic_subr4  r5  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 r8  )r  r  r(  r)  r   r:  r*  r   r+  r   r.  r/  r   r#   r1  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 r8  )r  r  r(  r)  r   r:  r*  r   r+  r   r.  r/  r   r#   r1  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 rC  )r  r  r(  r)  r   r*  rD  r   r+  r   r,  r-  r
  r.  r/  )r   r#   r1  cuda_atomic_sub_floatr4  r5  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 rI  )r  r  r(  r)  r   r:  r*  r   r+  r   r.  r/  rz  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 rI  )r  r  r(  r)  r   r:  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 rZ  )r  r  r(  r]  r,  r   r   r+  r   r-  r
  r^  r_  )r   r$   r#   rj  r4  r5  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 rd  )r  r  r(  r)  r   r:  r*  r   r+  r   r^  r_  ri  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 rd  )r  r  r(  r)  r   r:  r*  r   r+  r   r^  r_  ri  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 rZ  )r  r  r(  r]  r,  r   r   r+  r   r-  r
  r^  r_  )r   r$   r#   rn  rj  r4  r5  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 rd  )r  r  r(  r)  r   r:  r*  r   r+  r   r^  r_  ri  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 rd  )r  r  r(  r)  r   r:  r*  r   r+  r   r^  r_  ri  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   rB   r$  void(uint32[:], uint32)r%  )r  r  r(  r)  r   r*  r   r+  r   r-  r
  r.  r/  r   
rand_constr#   r1  rj  r4  r5  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   rB   r$  rN   rO   void(uint32[:,:], uint32)r9  )r  r  r(  r)  r   r:  r*  r   r+  r   r.  r/  r   r  r#   r1  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   r:  r*  r   r+  r   r.  r/  r   r  r#   r1  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   rB   r[  zvoid(int32[:], int32[:], int32)r%  )r  r  r(  int32r   r+  r   r*  r-  r
  r^  r_  r   r  r$   r#   rn  rj  r4  r5  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   r:  r*  r   r+  r   r^  r_  r   r  r#   r1  rj  s        r   test_atomic_and_global_2z(TestCudaAtomics.test_atomic_and_global_2.      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   rB   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_or6  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   r:  r*  r   r+  r   r.  r/  r  s        r   test_atomic_or2zTestCudaAtomics.test_atomic_or2C      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   r:  r*  r   r+  r   r.  r/  r  s        r   test_atomic_or3zTestCudaAtomics.test_atomic_or3K  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
  r^  r_  r  s           r   test_atomic_or_globalz%TestCudaAtomics.test_atomic_or_globalS  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   r:  r*  r   r+  r   r^  r_  r  s        r   test_atomic_or_global_2z'TestCudaAtomics.test_atomic_or_global_2a  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_xori  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   r:  r*  r   r+  r   r.  r/  )r   r  r#   r1  cuda_atomic_xor2s        r   test_atomic_xor2z TestCudaAtomics.test_atomic_xor2v  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   r:  r*  r   r+  r   r.  r/  )r   r  r#   r1  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
  r^  r_  )r   r  r$   r#   r4  rn  rj  r5  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   r:  r*  r   r+  r   r^  r_  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 )NrB   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 )NrB   r&  r   r$  rN   rO   )r  r  r(  r)  r:  )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 r3   r*  r   r+  r  r^  r_  where
r   r#   r$   r  rn  nblocksblksizer   r1  rj  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 r3   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 r3   r  	r   r#   r  rn  r  r  r   r1  rj  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)r4   rB   )r  r  r   r  r   r   r  r#   r$   rn  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)r4   rB   )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  r4   rM   )r  r  r   r  r   r   r  r#   rn  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)r4   rM   )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 r3   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 r3   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 r3   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_64-  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   rB   r$  r  r%  )r  r  r(  r   r)  r  r   r+  r   r^  r_  )r   r  r#   r$   rj  s        r   test_atomic_exchz TestCudaAtomics.test_atomic_exch2  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   rB   r$  rN   rO   r  r9  )r  r  r(  r   r)  r:  r   r+  r   r^  r_  r   r  r#   rj  s       r   test_atomic_exch2z!TestCudaAtomics.test_atomic_exch2<      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   rB   r$  rN   rO   r  r9  )r  r  r(  r   r)  r:  r   r+  r   r^  r_  r  s       r   test_atomic_exch3z!TestCudaAtomics.test_atomic_exch3D  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&  rB   r   r[  r  r%  )
r  r  r(  r   r  r   r+  r   r^  r_  )r   r  r$   r#   rn  rj  s         r   test_atomic_exch_globalz'TestCudaAtomics.test_atomic_exch_globalL  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rB   rB   r$  r4   r&  )r  r  r(  r)  r,  r'  r   r+  
atomic_maxmaxr^  r_  r   r'  lohivalsr  rj  r4  s           r   check_atomic_maxz TestCudaAtomics.check_atomic_maxV  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)  )r+  r  r  r   s    r   test_atomic_max_int32z%TestCudaAtomics.test_atomic_max_int32^  %    BHEBBBBBr   c                 J    |                      t          j        dd           d S Nr   r/  r0  )r+  r  r   r1  s    r   test_atomic_max_uint32z&TestCudaAtomics.test_atomic_max_uint32a  %    BI!>>>>>r   c                 J    |                      t          j        dd           d S r-  )r+  r  r]  r1  s    r   test_atomic_max_int64z%TestCudaAtomics.test_atomic_max_int64d  r3  r   c                 J    |                      t          j        dd           d S r5  )r+  r  r   r1  s    r   test_atomic_max_uint64z&TestCudaAtomics.test_atomic_max_uint64g  r7  r   c                 J    |                      t          j        dd           d S r-  )r+  r  r   r1  s    r   test_atomic_max_float32z'TestCudaAtomics.test_atomic_max_float32j  %    BJ6eDDDDDr   c                 J    |                      t          j        dd           d S r-  )r+  r  r   r1  s    r   test_atomic_max_doublez&TestCudaAtomics.test_atomic_max_doublem  r>  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   r/  r$  r$  r4   void(float64[:], float64[:,:]))r  r  r(  r)  r   r,  r   r+  !atomic_max_double_normalizedindexr&  r^  r_  r   r*  r  rj  r4  s        r   &test_atomic_max_double_normalizedindexz6TestCudaAtomics.test_atomic_max_double_normalizedindexp  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      rB   r$  r4   void(float64[:], float64[:])r%  )r  r  r(  r)  r   r,  r   r+  atomic_max_double_oneindexr&  r^  r_  rE  s        r   test_atomic_max_double_oneindexz/TestCudaAtomics.test_atomic_max_double_oneindexz  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$  r/  r&  )r  r  r(  r)  r!   r'  r   r+  
atomic_minminr^  r_  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 r-  )rP  r  r  r1  s    r   test_atomic_min_int32z%TestCudaAtomics.test_atomic_min_int32  r3  r   c                 J    |                      t          j        dd           d S r5  )rP  r  r   r1  s    r   test_atomic_min_uint32z&TestCudaAtomics.test_atomic_min_uint32  r7  r   c                 J    |                      t          j        dd           d S r-  )rP  r  r]  r1  s    r   test_atomic_min_int64z%TestCudaAtomics.test_atomic_min_int64  r3  r   c                 J    |                      t          j        dd           d S r5  )rP  r  r   r1  s    r   test_atomic_min_uint64z&TestCudaAtomics.test_atomic_min_uint64  r7  r   c                 J    |                      t          j        dd           d S r-  )rP  r  r   r1  s    r   test_atomic_min_floatz%TestCudaAtomics.test_atomic_min_float  r>  r   c                 J    |                      t          j        dd           d S r-  )rP  r  r   r1  s    r   test_atomic_min_doublez&TestCudaAtomics.test_atomic_min_double  r>  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 rB  )r  r  r(  r)  r   onesr   r+  !atomic_min_double_normalizedindexrO  r^  r_  rE  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 rH  )r  r  r(  r)  r   r^  r   r+  atomic_min_double_oneindexrO  r^  r_  rE  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 )NrC  r   rI  r4   r4   r$  r4   )r   r+  r  r  r(  r)  r   r,  nanr^  r_  )r   r   rj  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 )NrC  r   rI  r4   r$  re  )r   r+  r  r  r(  r)  r   r*  r,  rf  r^  r_  )r   r   rj  r  r4  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   )rg  rN  r1  s    r   test_atomic_min_nan_locationz,TestCudaAtomics.test_atomic_min_nan_location      --j99999r   c                 :    |                      t                     d S r   )rg  r%  r1  s    r   test_atomic_max_nan_locationz,TestCudaAtomics.test_atomic_max_nan_location  rl  r   c                 :    |                      t                     d S r   )ri  rN  r1  s    r   test_atomic_min_nan_valz'TestCudaAtomics.test_atomic_min_nan_val      ((44444r   c                 :    |                      t                     d S r   )ri  r%  r1  s    r   test_atomic_max_nan_valz'TestCudaAtomics.test_atomic_max_nan_val  rq  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   rB   r$  r4   rJ  r%  )r  r  r(  r)  r   r,  r   r+  atomic_max_double_sharedr&  r^  r_  r   r*  r  rn  rj  r4  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 ru  )r  r  r(  r)  r   r^  r   r+  atomic_min_double_sharedrO  r^  r_  rw  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   r4   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&  )
   r4   r}  r$  r}  r}  )r  r  )r  r  shuffleasarrayr<   
zeros_liker(  r)  r'  r*  r   r+  r^  assert_array_equal)r   nfillunfillr'  cas_funcndimr  outr#   	fill_maskunfill_mask
expect_res
expect_outrj  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  r1  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  r]  r  r1  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&  r4      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  r1  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  r1  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  r]  r  r1  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  r]  r  r1  s    r   test_atomic_cas2_2dimz%TestCudaAtomics.test_atomic_cas2_2dim,  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_1dim0  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&  r4   r  r  r  r  )r  r  r(  r   r  r  r  s      r   test_atomic_cas3_2dimz%TestCudaAtomics.test_atomic_cas3_2dim6  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_1dim<  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_2dimB  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   re  r4   )r  r,  r   isnanr.  assertEqualr   rV  initialr   s       r   _test_atomic_returns_oldz(TestCudaAtomics._test_atomic_returns_oldM  s    HQbj)))!tQ8G 	,OOBHQqTNN+++++QqT7+++++r   c                 \    t           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S r3   )r   rD   rE   r   s    r   rV  z;TestCudaAtomics.test_atomic_add_returns_old.<locals>.kernelW  !    ;??1a++AaDDDr   r}  r   r+  r  r   rV  s     r   test_atomic_add_returns_oldz+TestCudaAtomics.test_atomic_add_returns_oldV  ;    		, 	, 
	, 	%%fb11111r   c                 \    t           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S r3   r   rD   r&  r  s    r   rV  zBTestCudaAtomics.test_atomic_max_returns_no_replace.<locals>.kernel^  r  r   r}  r  r  s     r   "test_atomic_max_returns_no_replacez2TestCudaAtomics.test_atomic_max_returns_no_replace]  r  r   c                 \    t           j        d             }|                     |d           d S )Nc                 N    t           j                            | dd          | d<   d S Nr   r}  r4   r  r  s    r   rV  zCTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernele  !    ;??1a,,AaDDDr   r4   r  r  s     r   #test_atomic_max_returns_old_replacez3TestCudaAtomics.test_atomic_max_returns_old_replaced  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 r3   r  r  s    r   rV  zHTestCudaAtomics.test_atomic_max_returns_old_nan_in_array.<locals>.kernell  r  r   r   r+  r  r  rf  r  s     r   (test_atomic_max_returns_old_nan_in_arrayz8TestCudaAtomics.test_atomic_max_returns_old_nan_in_arrayk  s=    		, 	, 
	, 	%%fbf55555r   c                 \    t           j        d             }|                     |d           d S )Nc                 b    t           j                            | dt          j                  | d<   d S r3   )r   rD   r&  r  rf  r  s    r   rV  zCTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernels  #    ;??1a00AaDDDr   r}  r  r  s     r   #test_atomic_max_returns_old_nan_valz3TestCudaAtomics.test_atomic_max_returns_old_nan_valr  ;    		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      r4   r   rD   rO  r  s    r   rV  zFTestCudaAtomics.test_atomic_min_returns_old_no_replace.<locals>.kernelz  r  r   r}  r  r  s     r   &test_atomic_min_returns_old_no_replacez6TestCudaAtomics.test_atomic_min_returns_old_no_replacey  ;    		- 	- 
	- 	%%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   rV  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   rV  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 r3   )r   rD   rO  r  rf  r  s    r   rV  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$  r4   r  r&  )r  r  r(  r)  r,  r'  r   r+  atomic_nanmaxnanmaxr^  r_  	r   r'  r(  r)  init_valr*  r  rj  r4  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 Nr.  r/  r   r'  r(  r)  r  )r  r  r  r1  s    r   test_atomic_nanmax_int32z(TestCudaAtomics.test_atomic_nanmax_int32  4      rxFu*+ 	! 	- 	- 	- 	- 	-r   c                 L    |                      t          j        ddd           d S Nr   r/  r  )r  r  r   r1  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  r]  r1  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   r1  s    r   test_atomic_nanmax_uint64z)TestCudaAtomics.test_atomic_nanmax_uint64  r  r   c                 `    |                      t          j        ddt          j                   d S Nr.  r/  r  )r  r  r   rf  r1  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   rf  r1  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   rB   r$  r4   r  r&  rJ  r%  )r  r  r(  r)  r   rf  r!   r'  r   r+  atomic_nanmax_double_sharedr  r^  r_  rw  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   rI  rB   r$  r4   r  rJ  r%  )r  r  r(  r)  r   rf  r,  r   r+  rK  r  r^  r_  rE  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$  r4   r  r/  r&  )r  r  r(  r)  r!   r'  r   r+  atomic_nanminnanminr^  r_  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  r1  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   r1  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  r]  r1  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   r1  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   rf  r1  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   rf  r1  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   rf  r!   r'  r   r+  atomic_nanmin_double_sharedr  r^  r_  rw  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   rf  r!   r   r+  rb  r  r^  r_  rE  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   r4   re  )r  r,  r   rf  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 r3   r   rD   r  r  s    r   rV  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   rV  zFTestCudaAtomics.test_atomic_nanmax_returns_old_replace.<locals>.kernel  #    ;%%aB//AaDDDr   r4   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 r3   r#  r  s    r   rV  zKTestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array.<locals>.kernel$  r$  r   r   r+  r   r  rf  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 r3   )r   rD   r  r  rf  r  s    r   rV  zFTestCudaAtomics.test_atomic_nanmax_returns_old_nan_val.<locals>.kernel+  %    ;%%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   rD   r  r  s    r   rV  zITestCudaAtomics.test_atomic_nanmin_returns_old_no_replace.<locals>.kernel2  r)  r   r}  r%  r  s     r   )test_atomic_nanmin_returns_old_no_replacez9TestCudaAtomics.test_atomic_nanmin_returns_old_no_replace1  ;    		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  r6  r  s    r   rV  zFTestCudaAtomics.test_atomic_nanmin_returns_old_replace.<locals>.kernel9  r)  r   r  r%  r  s     r   &test_atomic_nanmin_returns_old_replacez6TestCudaAtomics.test_atomic_nanmin_returns_old_replace8  r8  r   c                 p    t           j        d             }|                     |t          j                   d S )Nc                 N    t           j                            | dd          | d<   d S r  r6  r  s    r   rV  zKTestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_array.<locals>.kernel@  r)  r   r-  r  s     r   +test_atomic_nanmin_returns_old_nan_in_arrayz;TestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_array?  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 r3   )r   rD   r  r  rf  r  s    r   rV  zFTestCudaAtomics.test_atomic_nanmin_returns_old_nan_val.<locals>.kernelG  r1  r   r  r%  r  s     r   &test_atomic_nanmin_returns_old_nan_valz6TestCudaAtomics.test_atomic_nanmin_returns_old_nan_valF  r3  r   )T)r4   )__name__
__module____qualname__r  r6  r=  rA  rG  rK  rM  rX  r   rb  rg  rk  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
  r  r  r  r  r  r  r  r  r   r"  r+  r2  r6  r9  r;  r=  r@  rF  rL  rP  rR  rT  rV  rX  rZ  r\  r`  rc  rg  ri  rk  rn  rp  rs  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.  r2  r7  r;  r>  rA  __classcell__)r!  s   @r   r  r    sG	           2 2 2$6 6 61 1 12 2 2$6 6 61 1 13 3 3 3" 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__)rnumpyr  textwrapr   numbar   r   r   r   r   numba.cuda.testingr	   r
   r   r   
numba.corer   r+  r   r   r   r/   r1   r:   r=   r?   rG   rJ   rQ   rT   rV   r]   ra   re   ri   rl   rr   ru   ry   r{   r}   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  r%  rD  rK  rv  rN  r_  rb  rz  r  $atomic_nanmax_double_normalizedindexatomic_nanmax_double_oneindexr  r
  $atomic_nanmin_double_normalizedindexatomic_nanmin_double_oneindexr  r  r  r  r  rB  mainr   r   r   <module>rQ     s             8 8 8 8 8 8 8 8 8 8 8 8 8 8/ / / / / / / / / / / /       
   
   
   
   
	 	 	 
   
   
  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u6 u6 u6 u6 u6l u6 u6 u6p# zHMOOOOO r   