+
    :iX                     6   ^ RI t^ RIHt ^ RIHtHtHtHtH	t	 ^ RI
HtHtHt ^ RIHt ]P                   ! RR7      R 4       t]P                   ! RR7      R	 4       t]P                   ! RR7      R
 4       t]P                   ! RR7      R 4       t]P                   ! RR7      R 4       t]P                   ! RR7      R 4       t]P                   ! RR7      R 4       t]P                   ! RR7      R 4       tR tR tR tR tR tR tR tR t R t!R t"R t#R t$R t%R t&R t'R t(R  t)R! t*R" t+R# t,R$ t-R% t.R& t/R' t0R( t1R) t2R* t3R+ t4R, t5R- t6R. t7R/ t8R0 t9R1 t:R2 t;R3 t<R4 t=R5 t>R6 t?R7 t@R8 tAR9 tBR: tCR; tDR< tER= tFR> tGR? tHR@ tIRA tJRB tKRC tLRD tMRE tNRF tORG tPRH tQRI tRRJ tSRK tTRL tURM tVRN tWRO tXRP tYRQ tZ]Z! RR4      w  t[t\t]t^]Z! RS4      w  t_t`tatb]Z! RT4      w  tctdtetf]Z! RU4      w  tgthtitjRV tkRW tlRX tm ! RY RZ]4      tn]oR[8X  d   ]P                  ! 4        R# R# )\    N)dedent)cudauint32uint64float32float64)unittestCUDATestCasecc_X_or_above)configT)devicec                     \        V 4      # N)r   nums   &b/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/tests/cudapy/test_atomics.pyatomic_cast_to_uint64r   	   s    #;    c                     \        V 4      # r   )intr   s   &r   atomic_cast_to_intr      s    s8Or   c                     V # r    r   s   &r   atomic_cast_noner      s    Jr   c	                 D   \         P                  P                  p	\         P                  P	                  WC4      p
WzV	&   \         P
                  ! 4        V! W,          V,          4      pV'       d	   W,          pV! WV4       \         P
                  ! 4        W,          W	&   R # 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.      su     ..

C			=	4BsG
CH},
-C!rwCHr   c                 0   \         P                  P                  p\         P                  P	                  WC4      pW,          W&   \         P
                  ! 4        V! W,          V,          4      p	V! WV4       \         P
                  ! 4        W,          W&   R # r   r   )
r"   r#   r$   r%   r&   r'   r(   r+   r,   r-   s
   &&&&&&&   r   atomic_binary_1dim_shared2r0   (   sl     ..

C			=	4BhBG
CH},
-CrwCHr   c                    \         P                  P                  p\         P                  P                  p\         P                  P                  W24      p	WV3,          WV3&   \         P                  ! 4        Wu! V4      3p
V'       d/   V
^ ,          V^ ,          ,          V
^,          V^,          ,          3p
V! WV4       \         P                  ! 4        WV3,          WV3&   R# r   N)r   r   r   yr   r    r!   )r"   r$   r%   	ary_shaper'   y_cast_funcr*   txtyr,   r-   s   &&&&&&&    r   atomic_binary_2dim_sharedr8   5   s     
		B			B			9	0BVB2vJ{2
C1v	!$c!fy|&;<rV*CBKr   c                 &   \         P                  P                  p\         P                  P                  pWS! V4      3pV'       dC   V^ ,          V P                  ^ ,          ,          V^,          V P                  ^,          ,          3pV! WV4       R# r2   )r   r   r   r3   shape)r"   r$   r'   r5   r*   r6   r7   r-   s   &&&&&   r   atomic_binary_2dim_globalr;   E   sg    			B			B{2
C1v		!$c!fsyy|&;<sr   c                     \         P                  P                  p\        W,          V,          4      pV'       d	   Wr,          pV! WV4       R # r   )r   r   r   r   )r"   r#   r&   r$   r'   r*   r+   r-   s   &&&&&&  r   atomic_binary_1dim_globalr=   O   s9     ..

C
ch&
'C!sr   c                 j    \        W ^\        ^ \        P                  P                  \
        ^ R4	       R#    FNr.   r   r   atomicaddr   r"   s   &r   
atomic_addrE   Y   %    c62"kkoo/?EKr   c                 j    \        W ^\        ^ \        P                  P                  \
        ^ R4	       R# r@   TNrA   rD   s   &r   atomic_add_wraprI   ^   s%    c62"kkoo/?DJr   c           	      h    \        V ^\        R\        P                  P                  \
        R4       R# r@   FN      r8   r   r   rB   rC   r   rD   s   &r   atomic_add2rP   c   #    c1ff"kkoo/?Hr   c           	      h    \        V ^\        R\        P                  P                  \
        R4       R# )r@   TNrL   rO   rD   s   &r   atomic_add2_wraprS   h   s#    c1ff"kkoo/?Gr   c           	      h    \        V ^\        R\        P                  P                  \
        R4       R# rK   )r8   r   r   rB   rC   r   rD   s   &r   atomic_add3rU   m   #    c1ff"kkoo/DeMr   c                 j    \        W R \        ^ \        P                  P                  \
        RR4	       R#       ?        FNr.   r   r   rB   rC   r   rD   s   &r   atomic_add_floatr\   r   %    cWb"kkoo/A3Or   c                 j    \        W R \        ^ \        P                  P                  \
        RR4	       R# rY   rZ   TNr[   rD   s   &r   atomic_add_float_wrapr`   w   s%    cWb"kkoo/A3Nr   c           	      h    \        V R \        R\        P                  P                  \
        R4       R# rY   FNrL   r8   r   r   rB   rC   r   rD   s   &r   atomic_add_float_2rd   |   #    c3"kkoo/?Hr   c           	      h    \        V R \        R\        P                  P                  \
        R4       R# rY   TNrL   rc   rD   s   &r   atomic_add_float_2_wraprh      #    c3"kkoo/?Gr   c           	      h    \        V R \        R\        P                  P                  \
        R4       R# rb   )r8   r   r   rB   rC   r   rD   s   &r   atomic_add_float_3rk      #    c3"kkoo/DeMr   c                 T    \        W^ R\        P                  P                  R4       R#     rY   FNr=   r   rB   rC   r#   r"   s   &&r   atomic_add_double_globalrr          cC%Hr   c                 T    \        W^ R\        P                  P                  R4       R# )ro   rY   TNrp   rq   s   &&r   atomic_add_double_global_wrapru      s    cC$Gr   c                 \    \        V ^\        P                  P                  \        R4       R# r?   r;   r   rB   rC   r   rD   s   &r   atomic_add_double_global_2rx      s    c1dkkoo7GOr   c                 \    \        V ^\        P                  P                  \        R4       R# rH   rw   rD   s   &r   atomic_add_double_global_2_wraprz      s    c1dkkoo7GNr   c                 \    \        V ^\        P                  P                  \        R4       R# r?   )r;   r   rB   rC   r   rD   s   &r   atomic_add_double_global_3r|      s    c1dkkoo7L#%r   c                 j    \        WR \        ^ \        P                  P                  \
        RR4	       R# rX   r.   r   r   rB   rC   r   rq   s   &&r   atomic_add_doubler      %    cWb"kkoo/?eMr   c                 j    \        WR \        ^ \        P                  P                  \
        RR4	       R# r_   r~   rq   s   &&r   atomic_add_double_wrapr      s%    cWb"kkoo/?dLr   c           	      h    \        V R \        R\        P                  P                  \
        R4       R# rb   r8   r   r   rB   rC   r   rD   s   &r   atomic_add_double_2r      re   r   c           	      h    \        V R \        R\        P                  P                  \
        R4       R# rg   r   rD   s   &r   atomic_add_double_2_wrapr      ri   r   c           	      h    \        V R \        R\        P                  P                  \
        R4       R# rb   )r8   r   r   rB   rC   r   rD   s   &r   atomic_add_double_3r      rl   r   c                 j    \        W ^\        ^ \        P                  P                  \
        ^ R4	       R# r?   )r.   r   r   rB   subr   rD   s   &r   
atomic_subr      rF   r   c           	      h    \        V ^\        R\        P                  P                  \
        R4       R# rK   )r8   r   r   rB   r   r   rD   s   &r   atomic_sub2r      rQ   r   c           	      h    \        V ^\        R\        P                  P                  \
        R4       R# rK   )r8   r   r   rB   r   r   rD   s   &r   atomic_sub3r      rV   r   c                 j    \        W R \        ^ \        P                  P                  \
        RR4	       R# rX   )r.   r   r   rB   r   r   rD   s   &r   atomic_sub_floatr      r]   r   c           	      h    \        V R \        R\        P                  P                  \
        R4       R# rb   )r8   r   r   rB   r   r   rD   s   &r   atomic_sub_float_2r      re   r   c           	      h    \        V R \        R\        P                  P                  \
        R4       R# rb   )r8   r   r   rB   r   r   rD   s   &r   atomic_sub_float_3r      rl   r   c                 j    \        WR \        ^ \        P                  P                  \
        RR4	       R# rX   )r.   r   r   rB   r   r   rq   s   &&r   atomic_sub_doubler      r   r   c           	      h    \        V R \        R\        P                  P                  \
        R4       R# rb   )r8   r   r   rB   r   r   rD   s   &r   atomic_sub_double_2r      re   r   c           	      h    \        V R \        R\        P                  P                  \
        R4       R# rb   r8   r   r   rB   r   r   rD   s   &r   atomic_sub_double_3r      rl   r   c                 T    \        W^ R\        P                  P                  R4       R# rn   )r=   r   rB   r   rq   s   &&r   atomic_sub_double_globalr      rs   r   c                 \    \        V R \        P                  P                  \        R4       R# )rY   FN)r;   r   rB   r   r   rD   s   &r   atomic_sub_double_global_2r      s    c39I#%r   c           	      h    \        V R \        R\        P                  P                  \
        R4       R# rb   r   rD   s   &r   atomic_sub_double_global_3r      rl   r   c                 j    \        W V\        ^ \        P                  P                  \
        ^R4	       R# ro   FN)r.   r   r   rB   and_r   r"   r$   s   &&r   
atomic_andr      s'    cVR"kk..0@!ULr   c           	      f    \        W\        R\        P                  P                  \
        R4       R# rM   FNrL   )r8   r   r   rB   r   r   r   s   &&r   atomic_and2r      #    c"kk..0@%Ir   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   atomic_and3r      s#    c"kk..0EuNr   c                 T    \        W^ V\        P                  P                  R4       R# r   )r=   r   rB   r   r#   r"   r$   s   &&&r   atomic_and_globalr         cC1A1A5Ir   c                 Z    \        W\        P                  P                  \        R 4       R# FN)r;   r   rB   r   r   r   s   &&r   atomic_and_global_2r     s    c(8(8.7r   c                 j    \        W V\        ^ \        P                  P                  \
        ^ R4	       R# r   )r.   r   r   rB   or_r   r   s   &&r   	atomic_orr     %    cVR"kkoo/?EKr   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   
atomic_or2r     !    c"kkoo/?Hr   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   
atomic_or3r     !    c"kkoo/DeMr   c                 T    \        W^ V\        P                  P                  R4       R# r   )r=   r   rB   r   r   s   &&&r   atomic_or_globalr     rs   r   c                 Z    \        W\        P                  P                  \        R 4       R# r   )r;   r   rB   r   r   r   s   &&r   atomic_or_global_2r         c.7r   c                 j    \        W V\        ^ \        P                  P                  \
        ^ R4	       R# r   )r.   r   r   rB   xorr   r   s   &&r   
atomic_xorr   $  r   r   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   atomic_xor2r   )  r   r   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   atomic_xor3r   .  r   r   c                 T    \        W^ V\        P                  P                  R4       R# r   )r=   r   rB   r   r   s   &&&r   atomic_xor_globalr   3  rs   r   c                 Z    \        W\        P                  P                  \        R 4       R# r   )r;   r   rB   r   r   r   s   &&r   atomic_xor_global_2r   7  r   r   c           	      f    \        WV\        ^ \        P                  P                  \
        4       R# ro   N)r0   r   r   rB   incr   r"   r#   r$   s   &&&r   atomic_inc32r   <  !    sfb#{{0@Br   c           	      f    \        WV\        ^ \        P                  P                  \
        4       R# r   )r0   r   r   rB   r   r   r   s   &&&r   atomic_inc64r   A  !    sfb#{{0BDr   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   atomic_inc2_32r   F  r   r   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   atomic_inc2_64r   K  r   r   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   atomic_inc3r   P  r   r   c                 T    \        W^ V\        P                  P                  R4       R# r   )r=   r   rB   r   r   s   &&&r   atomic_inc_globalr   U  rs   r   c                 Z    \        W\        P                  P                  \        R 4       R# r   )r;   r   rB   r   r   r   s   &&r   atomic_inc_global_2r   Y  r   r   c           	      f    \        WV\        ^ \        P                  P                  \
        4       R# r   )r0   r   r   rB   decr   r   s   &&&r   atomic_dec32r   ^  r   r   c           	      f    \        WV\        ^ \        P                  P                  \
        4       R# r   )r0   r   r   rB   r   r   r   s   &&&r   atomic_dec64r   c  r   r   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   atomic_dec2_32r   h  r   r   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   atomic_dec2_64r   m  r   r   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   atomic_dec3r   r  r   r   c                 T    \        W^ V\        P                  P                  R4       R# r   )r=   r   rB   r   r   s   &&&r   atomic_dec_globalr   w  rs   r   c                 Z    \        W\        P                  P                  \        R 4       R# r   )r;   r   rB   r   r   r   s   &&r   atomic_dec_global_2r   {  r   r   c           	      f    \        WV\        ^ \        P                  P                  \
        4       R# r   )r0   r   r   rB   exchr   r   s   &&&r   atomic_exchr     s#    sfb#{{//1ACr   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   atomic_exch2r     r   r   c           	      f    \        W\        R\        P                  P                  \
        R4       R# r   )r8   r   r   rB   r   r   r   s   &&r   atomic_exch3r     r   r   c                 T    \        W^ V\        P                  P                  R4       R# r   )r=   r   rB   r   r   s   &&&r   atomic_exch_globalr     r   r   c           	          \        R 4      P                  V R7      p/ p\        VR\        R\        R\
        /V4       VR,          VR,          VR,          VR,          3# )	a  
    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]
    )funcr   r   r   rB   atomic_double_normalizedindexatomic_double_oneindexatomic_double_shared)r   formatexecr   r   r   )r   fnslds   &  r   gen_atomic_extreme_funcsr    sn    
  	6 
T	7 8 
BvtY6BBGxL"<='("-C*DF Fr   zcuda.atomic.maxzcuda.atomic.minzcuda.atomic.nanmaxzcuda.atomic.nanminc                     \         P                  ! ^4      pW@P                  8  d.   \         P                  P	                  WR W2V,          4      W&   R# R# r@   N)r   gridsizerB   compare_and_swapresoldr"   fill_valgids   &&&& r   atomic_compare_and_swapr    s=    
))A,C
XX~;;//D	8XN r   c                     \         P                  ! ^4      pW@P                  8  d,   \         P                  P	                  WW2V,          4      W&   R# R# r  )r   r  r  rB   casr
  s   &&&& r   atomic_cas_1dimr    s7    
))A,C
XX~;;??3X3x@ r   c                    \         P                  ! ^4      pV^ ,          V P                  ^ ,          8  dM   V^,          V P                  ^,          8  d,   \         P                  P	                  WW2V,          4      W&   R# R# R# )   N)r   r  r:   rB   r  r
  s   &&&& r   atomic_cas_2dimr    sX    
))A,C
1v		!Q#))A,!6;;??3X3x@ "7r   c                     a a ] tR tRt oV 3R ltR tR tR tR tR t	R t
RR	 ltR
 tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR t R t!R  t"R! t#R" t$R# t%R$ t&R% t'R& t(R' t)R( t*R) t+R* t,R+ t-R, t.R- t/R. t0R/ t1R0 t2R1 t3R2 t4R3 t5R4 t6R5 t7R6 t8R7 t9R8 t:R9 t;R: t<R; t=R< t>R= t?R> t@R? tAR@ tBRA tCRB tDRC tERD tFRE tGRF tHRG tIRH tJRI tKRJ tLRK tMRL tNRM tORN tPRO tQRP tRRQ tSRR tTRS tURT tVRU tWRV tXRW tYRX tZRY t[RZ t\R[ t]R\ t^R] t_R^ t`R_ taR` tbRa tcRb tdRRc lteRd tfRe tgRf thRg tiRh tjRi tkRj tlRk tmRl tnRm toRn tpRo tqRp trRq tsRr ttRs tuRt tvRu twRv txRw tyRx tzRy t{Rz t|R{ t}R| t~R} tR~ tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tRtVtV ;t# )TestCudaAtomicsi  c                b   < \         SV `  4        \        P                  P	                  ^ 4       R# r2   )supersetUpnprandomseed)self	__class__s   &r   r  TestCudaAtomics.setUp  s    
		qr   c                   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      pVP                  4       pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       \        P                  ! R4      ! \        4      pVR,          ! V4       \         P                  ! ^ \         P                  R7      p\        VP                  4       F  pWcV,          ;;,          ^,          uu&   K   	  V P                  \         P                  ! W8H  4      4       V P                  \         P                  ! W&8H  4      4       R# r   r  zvoid(uint32[:])dtypeNr@   ro   )r  r  randintastyper   copyr   jitrE   rI   zerosranger  
assertTrueall)r  r"   ary_wraporigcuda_atomic_addcuda_atomic_add_wrapgoldis   &       r   test_atomic_addTestCudaAtomics.test_atomic_add  s    ii2B/66ryyA88:xxz((#45jAs##xx(9:?KU#H-xx")),tyy!AaMQM " 	s{+,x/01r   c                X   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       \        P                  ! R4      ! \        4      pVR,          ! V4       V P                  \         P                  ! W^,           8H  4      4       V P                  \         P                  ! W#^,           8H  4      4       R# r   r#  zvoid(uint32[:,:])Nr@   rL   )r  r  r'  r(  r   reshaper)  r   r*  rP   rS   r-  r.  )r  r"   r/  r0  cuda_atomic_add2cuda_atomic_add2_wraps   &     r   test_atomic_add2 TestCudaAtomics.test_atomic_add2  s    ii2B/66ryyAII!QO88:xxz88$78E#C( $)< =>N Oi(2sQh/0x!8345r   c                ~   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       V P                  \         P                  ! W^,           8H  4      4       R# r8  )r  r  r'  r(  r   r:  r)  r   r*  rU   r-  r.  r  r"   r0  cuda_atomic_add3s   &   r   test_atomic_add3 TestCudaAtomics.test_atomic_add3  s    ii2B/66ryyAII!QOxxz88$78E#C(sQh/0r   c                   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      pVP                  4       pVP                  4       P                  \         P                  4      p\        P                  ! R4      ! \        4      pVR,          ! V4       \        P                  ! R4      ! \        4      pVR,          ! V4       \         P                  ! ^ \         P                  R7      p\        VP                  4       F  pWcV,          ;;,          R,          uu&   K   	  V P                  \         P                   ! W8H  4      4       V P                  \         P                   ! W&8H  4      4       R# r   r#  zvoid(float32[:])r$  rY   Nr&  )r  r  r'  r(  r   r)  intpr   r*  r\   r`   r+  r   r,  r  r-  r.  )r  r"   r/  r0  cuda_atomic_add_floatadd_float_wrapr3  r4  s   &       r   test_atomic_add_float%TestCudaAtomics.test_atomic_add_float  s    ii2B/66rzzB88:xxz  ) $); <=M Ne$S)"456KLuh'xx")),tyy!AaMS M " 	s{+,x/01r   c                X   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       \        P                  ! R4      ! \        4      pVR,          ! V4       V P                  \         P                  ! W^,           8H  4      4       V P                  \         P                  ! W#^,           8H  4      4       R# r   r#  zvoid(float32[:,:])Nr9  )r  r  r'  r(  r   r:  r)  r   r*  rd   rh   r-  r.  )r  r"   r/  r0  r;  cuda_func_wraps   &     r   test_atomic_add_float_2'TestCudaAtomics.test_atomic_add_float_2  s    ii2B/66rzzBJJ1aP88:xxz88$89:LM#C("678OPy!(+sQh/0x!8345r   c                ~   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       V P                  \         P                  ! W^,           8H  4      4       R# rL  )r  r  r'  r(  r   r:  r)  r   r*  rk   r-  r.  r@  s   &   r   test_atomic_add_float_3'TestCudaAtomics.test_atomic_add_float_3"  s    ii2B/66rzzBJJ1aPxxz88$89:LM#C(sQh/0r   c                   \         P                  '       d   R # \        \        VP	                  4       P                  4       4      4      p\        ^^ 4      '       dM   \        P                  P                  4       R8  d   RpMRpV'       d   V R2pV P                  V R2V4       R # V'       d   V P                  RV4       R # V P                  RV4       R # )Nredatomz.sharedz.add.f64zatom.shared.cas.b64zatom.cas.b64)   r@   )r   ENABLE_CUDASIMnextiterinspect_asmvaluesr   r   runtimeget_versionassertIn)r  kernelr   asminsts   &&&  r   assertCorrectFloat64Atomics+TestCudaAtomics.assertCorrectFloat64Atomics*  s        4**,33567A||'')G3 w'MMTF(+S13S9nc2r   c                   \         P                  P                  ^ ^ ^ \         P                  R7      p\         P                  ! ^ \         P
                  4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W4       \        P                  ! R4      ! \        4      pVR,          ! W4       \         P                  ! ^ \         P                  R7      p\        VP                  4       F  pWaV,          ;;,          R,          uu&   K   	  \         P                  P                  W&4       \         P                  P                  W64       V P!                  V4       V P!                  V4       R# r   r  r%  void(int64[:], float64[:])r$  rY   Nr&  )r  r  r'  int64r+  r   r)  r   r*  r   r   r   r,  r  testingassert_equalrb  )r  r#   r"   r/  cuda_fnwrap_fnr3  r4  s   &       r   test_atomic_add_double&TestCudaAtomics.test_atomic_add_doubleC  s    ii2Bbhh?hhr2::&88:((789JKs ((789OPs%xx")),sxxAQLCL ! 	

*


/((1((1r   c                |   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       \        P                  ! R4      ! \        4      pVR,          ! V4       \         P                  P                  W^,           4       \         P                  P                  W#^,           4       V P                  V4       V P                  V4       R# r   r#  void(float64[:,:])Nr9  )r  r  r'  r(  r   r:  r)  r   r*  r   r   ri  rj  rb  )r  r"   r/  r0  rk  cuda_fn_wraps   &     r   test_atomic_add_double_2(TestCudaAtomics.test_atomic_add_double_2W  s    ii2B/66rzzBJJ1aP88:xxz((/01DE	3xx 456NOY)


AX.


(3((1((6r   c                   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       \         P                  P                  W^,           4       V P                  V4       R# rp  )r  r  r'  r(  r   r:  r)  r   r*  r   ri  rj  rb  r  r"   r0  	cuda_funcs   &   r   test_atomic_add_double_3(TestCudaAtomics.test_atomic_add_double_3g  s    ii2B/66rzzBJJ1aPxxzHH123FG	)S!


AX.((3r   c                   \         P                  P                  ^ ^ ^ \         P                  R7      p\         P                  ! ^ \         P
                  4      pVP                  4       pRp\        P                  ! V4      ! \        4      p\        P                  ! V4      ! \        4      pVR,          ! W4       VR,          ! W4       \         P                  ! ^ \         P                  R7      p\        VP                  4       F  pWqV,          ;;,          R,          uu&   K   	  \         P                  P                  W'4       \         P                  P                  W74       V P!                  VRR7       V P!                  VRR7       R# )	r   rf  rg  r$  rY   Fr   Nr&  )r  r  r'  rh  r+  r   r)  r   r*  rr   ru   r   r,  r  ri  rj  rb  )	r  r#   r"   r/  sigrw  wrap_cuda_funcr3  r4  s	   &        r   test_atomic_add_double_global-TestCudaAtomics.test_atomic_add_double_globalp  s	   ii2Bbhh?hhr2::&88:*HHSM":;	#'DE%"uc,xx")),sxxAQLCL ! 	

*


/((5(A(((Fr   c                   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       pVP                  4       pRp\        P                  ! V4      ! \        4      p\        P                  ! V4      ! \        4      pVR,          ! V4       VR,          ! V4       \         P                  P                  W^,           4       \         P                  P                  W#^,           4       V P                  VRR7       V P                  VRR7       R# r   r#  rq  Fr{  Nr9  )r  r  r'  r(  r   r:  r)  r   r*  rx   rz   ri  rj  rb  )r  r"   r/  r0  r|  rw  r}  s   &      r   test_atomic_add_double_global_2/TestCudaAtomics.test_atomic_add_double_global_2  s    ii2B/66rzzBJJ1aP88:xxz"HHSM"<=	#'FG)S!y!(+


AX.


(3((5(A(((Fr   c                   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       \         P                  P                  W^,           4       V P                  VRR7       R# r  )r  r  r'  r(  r   r:  r)  r   r*  r|   ri  rj  rb  rv  s   &   r   test_atomic_add_double_global_3/TestCudaAtomics.test_atomic_add_double_global_3  s    ii2B/66rzzBJJ1aPxxzHH123MN	)S!


AX.((5(Ar   c                
   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       \         P                  ! ^ \         P                  R7      p\        VP                  4       F  pWBV,          ;;,          ^,          uu&   K   	  V P                  \         P                  ! W8H  4      4       R# r"  )r  r  r'  r(  r   r)  r   r*  r   r+  r,  r  r-  r.  )r  r"   r0  cuda_atomic_subr3  r4  s   &     r   test_atomic_subTestCudaAtomics.test_atomic_sub  s    ii2B/66ryyAxxz((#45jAs#xx")),tyy!AaMQM " 	s{+,r   c                ~   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       V P                  \         P                  ! W^,
          8H  4      4       R# r8  )r  r  r'  r(  r   r:  r)  r   r*  r   r-  r.  r  r"   r0  cuda_atomic_sub2s   &   r   test_atomic_sub2 TestCudaAtomics.test_atomic_sub2      ii2B/66ryyAII!QOxxz88$78E#C(sQh/0r   c                ~   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       V P                  \         P                  ! W^,
          8H  4      4       R# r8  )r  r  r'  r(  r   r:  r)  r   r*  r   r-  r.  r  r"   r0  cuda_atomic_sub3s   &   r   test_atomic_sub3 TestCudaAtomics.test_atomic_sub3  r  r   c                D   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      pVP                  4       P                  \         P                  4      p\        P                  ! R4      ! \        4      pVR,          ! V4       \         P                  ! ^ \         P                  R7      p\        VP                  4       F  pWBV,          ;;,          R,          uu&   K   	  V P                  \         P                  ! W8H  4      4       R# rE  )r  r  r'  r(  r   r)  rF  r   r*  r   r+  r,  r  r-  r.  )r  r"   r0  cuda_atomic_sub_floatr3  r4  s   &     r   test_atomic_sub_float%TestCudaAtomics.test_atomic_sub_float  s    ii2B/66rzzBxxz  ) $); <=M Ne$S)xx"**-tyy!AaMS M " 	s{+,r   c                ~   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       V P                  \         P                  ! W^,
          8H  4      4       R# rL  )r  r  r'  r(  r   r:  r)  r   r*  r   r-  r.  r  s   &   r   test_atomic_sub_float_2'TestCudaAtomics.test_atomic_sub_float_2      ii2B/66rzzBJJ1aPxxz88$89:LM#C(sQh/0r   c                ~   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       V P                  \         P                  ! W^,
          8H  4      4       R# rL  )r  r  r'  r(  r   r:  r)  r   r*  r   r-  r.  r  s   &   r   test_atomic_sub_float_3'TestCudaAtomics.test_atomic_sub_float_3  r  r   c                   \         P                  P                  ^ ^ ^ \         P                  R7      p\         P                  ! ^ \         P
                  4      p\        P                  ! R4      ! \        4      pVR,          ! W4       \         P                  ! ^ \         P
                  R7      p\        VP                  4       F  pWAV,          ;;,          R,          uu&   K   	  \         P                  P                  W$4       R# re  )r  r  r'  rh  r+  r   r   r*  r   r,  r  ri  rj  )r  r#   r"   rw  r3  r4  s   &     r   test_atomic_sub_double&TestCudaAtomics.test_atomic_sub_double  s    ii2Bbhh?hhr2::&HH9:;LM	%"xx"**-sxxAQLCL ! 	

*r   c                n   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       \         P                  P                  W^,
          4       R# rp  )r  r  r'  r(  r   r:  r)  r   r*  r   ri  rj  rv  s   &   r   test_atomic_sub_double_2(TestCudaAtomics.test_atomic_sub_double_2  |    ii2B/66rzzBJJ1aPxxzHH123FG	)S!


AX.r   c                n   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       \         P                  P                  W^,
          4       R# rp  )r  r  r'  r(  r   r:  r)  r   r*  r   ri  rj  rv  s   &   r   test_atomic_sub_double_3(TestCudaAtomics.test_atomic_sub_double_3  r  r   c                   \         P                  P                  ^ ^ ^ \         P                  R7      p\         P                  ! ^ \         P
                  4      pRp\        P                  ! V4      ! \        4      pVR,          ! W4       \         P                  ! ^ \         P
                  R7      p\        VP                  4       F  pWQV,          ;;,          R,          uu&   K   	  \         P                  P                  W%4       R# re  )r  r  r'  rh  r+  r   r   r*  r   r,  r  ri  rj  )r  r#   r"   r|  rw  r3  r4  s   &      r   test_atomic_sub_double_global-TestCudaAtomics.test_atomic_sub_double_global  s    ii2Bbhh?hhr2::&*HHSM":;	%"xx"**-sxxAQLCL ! 	

*r   c                n   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       \         P                  P                  W^,
          4       R# rp  )r  r  r'  r(  r   r:  r)  r   r*  r   ri  rj  rv  s   &   r   test_atomic_sub_double_global_2/TestCudaAtomics.test_atomic_sub_double_global_2  |    ii2B/66rzzBJJ1aPxxzHH123MN	)S!


AX.r   c                n   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! V4       \         P                  P                  W^,
          4       R# rp  )r  r  r'  r(  r   r:  r)  r   r*  r   ri  rj  rv  s   &   r   test_atomic_sub_double_global_3/TestCudaAtomics.test_atomic_sub_double_global_3  r  r   c                   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       VP                  4       p\        VP                  4       F  pWSV,          ;;,          V,          uu&   K   	  V P                  \         P                  ! W%8H  4      4       R# )  r#  void(uint32[:], uint32)Nr&  )r  r  r'  r(  r   r)  r   r*  r   r,  r  r-  r.  r  
rand_constr"   r0  rw  r3  r4  s   &      r   test_atomic_andTestCudaAtomics.test_atomic_and  s    YY&&s+
ii2B/66ryyAxxzHH67
C	%)xxztyy!AaMZ'M " 	s{+,r   c                   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       V P                  \         P                  ! W#V,          8H  4      4       R# r  r#  void(uint32[:,:], uint32)Nr9  )r  r  r'  r(  r   r:  r)  r   r*  r   r-  r.  r  r  r"   r0  cuda_atomic_and2s   &    r   test_atomic_and2 TestCudaAtomics.test_atomic_and2      YY&&s+
ii2B/66ryyAII!QOxxz88$?@M#C4sZ&7789r   c                   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       V P                  \         P                  ! W#V,          8H  4      4       R# r  )r  r  r'  r(  r   r:  r)  r   r*  r   r-  r.  r  r  r"   r0  cuda_atomic_and3s   &    r   test_atomic_and3 TestCudaAtomics.test_atomic_and3  r  r   c                8   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ \         P                  R7      p\         P                  P                  ^ ^ ^ \         P                  R7      pRp\        P
                  ! V4      ! \        4      pVR,          ! W#V4       VP                  4       p\        VP                  4       F  pWbV,          ;;,          V,          uu&   K   	  \         P                  P                  W64       R# r  rf  zvoid(int32[:], int32[:], int32)Nr&  )r  r  r'  int32r   r*  r   r)  r,  r  ri  rj  r  r  r#   r"   r|  rw  r3  r4  s   &       r   test_atomic_and_global&TestCudaAtomics.test_atomic_and_global%  s    YY&&s+
ii2Bbhh?ii2Bbhh?/HHSM"34	%:.xxzsxxAQLJ&L ! 	

*r   c                   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  P                  W#V,          4       R# r  )r  r  r'  r(  r   r:  r)  r   r*  r   ri  rj  r  r  r"   r0  rw  s   &    r   test_atomic_and_global_2(TestCudaAtomics.test_atomic_and_global_23      YY&&s+
ii2B/66ryyAII!QOxxzHH89:MN	)S-


J%67r   c                H   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  ! ^ \         P                  R7      p\        VP                  4       F  pWSV,          ;;,          V,          uu&   K   	  V P                  \         P                  ! W%8H  4      4       R# r  r#  r  r$  Nr&  )r  r  r'  r(  r   r)  r   r*  r   r+  r,  r  r-  r.  r  s   &      r   test_atomic_orTestCudaAtomics.test_atomic_or;  s    YY&&s+
ii2B/66ryyAxxzHH67	B	%)xx")),tyy!AaMZ'M " 	s{+,r   c                   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       V P                  \         P                  ! W#V,          8H  4      4       R# r  )r  r  r'  r(  r   r:  r)  r   r*  r   r-  r.  r  s   &    r   test_atomic_or2TestCudaAtomics.test_atomic_or2H      YY&&s+
ii2B/66ryyAII!QOxxz88$?@L#C4sZ&7789r   c                   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       V P                  \         P                  ! W#V,          8H  4      4       R# r  )r  r  r'  r(  r   r:  r)  r   r*  r   r-  r.  r  s   &    r   test_atomic_or3TestCudaAtomics.test_atomic_or3P  r  r   c                8   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ \         P                  R7      p\         P                  P                  ^ ^ ^ \         P                  R7      pRp\        P
                  ! V4      ! \        4      pVR,          ! W#V4       VP                  4       p\        VP                  4       F  pWbV,          ;;,          V,          uu&   K   	  \         P                  P                  W64       R# r  )r  r  r'  r  r   r*  r   r)  r,  r  ri  rj  r  s   &       r   test_atomic_or_global%TestCudaAtomics.test_atomic_or_globalX  s    YY&&s+
ii2Bbhh?ii2Bbhh?/HHSM"23	%:.xxzsxxAQLJ&L ! 	

*r   c                   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  P                  W#V,          4       R# r  )r  r  r'  r(  r   r:  r)  r   r*  r   ri  rj  r  s   &    r   test_atomic_or_global_2'TestCudaAtomics.test_atomic_or_global_2f  s    YY&&s+
ii2B/66ryyAII!QOxxzHH89:LM	)S-


J%67r   c                H   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  ! ^ \         P                  R7      p\        VP                  4       F  pWSV,          ;;,          V,          uu&   K   	  V P                  \         P                  ! W%8H  4      4       R# r  )r  r  r'  r(  r   r)  r   r*  r   r+  r,  r  r-  r.  r  s   &      r   test_atomic_xorTestCudaAtomics.test_atomic_xorn  s    YY&&s+
ii2B/66ryyAxxzHH67
C	%)xx")),tyy!AaMZ'M " 	s{+,r   c                   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       V P                  \         P                  ! W#V,          8H  4      4       R# r  )r  r  r'  r(  r   r:  r)  r   r*  r   r-  r.  )r  r  r"   r0  cuda_atomic_xor2s   &    r   test_atomic_xor2 TestCudaAtomics.test_atomic_xor2{  r  r   c                   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       V P                  \         P                  ! W#V,          8H  4      4       R# r  )r  r  r'  r(  r   r:  r)  r   r*  r   r-  r.  )r  r  r"   r0  cuda_atomic_xor3s   &    r   test_atomic_xor3 TestCudaAtomics.test_atomic_xor3  r  r   c                8   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ \         P                  R7      p\         P                  P                  ^ ^ ^ \         P                  R7      pVP	                  4       pRp\
        P                  ! V4      ! \        4      pVR,          ! W#V4       \        VP                  4       F  pWBV,          ;;,          V,          uu&   K   	  \         P                  P                  W44       R# r  )r  r  r'  r  r)  r   r*  r   r,  r  ri  rj  )r  r  r#   r"   r3  r|  rw  r4  s   &       r   test_atomic_xor_global&TestCudaAtomics.test_atomic_xor_global  s    YY&&s+
ii2Bbhh?ii2Bbhh?xxz/HHSM"34	%:.sxxAQLJ&L ! 	

*r   c                   \         P                  P                  R 4      p\         P                  P                  ^ ^ ^ R7      P                  \         P                  4      P                  ^^4      pVP                  4       p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  P                  W#V,          4       R# r  )r  r  r'  r(  r   r:  r)  r   r*  r   ri  rj  r  s   &    r   test_atomic_xor_global_2(TestCudaAtomics.test_atomic_xor_global_2  r  r   c                    \         P                  P                  ^ VR7      p\         P                  P                  ^ ^ ^ R7      P                  V4      p\         P                  ! ^ VR7      pW#V3# ro   r$  r#  )r  r  r'  r(  arange)r  r%  rconstraryary_idxs   &&   r   inc_dec_1dim_setup"TestCudaAtomics.inc_dec_1dim_setup  s\    ""2e"4yy  BR 077>))Be,W$$r   c                    \         P                  P                  ^ VR7      p\         P                  P                  ^ ^ ^ R7      P                  V4      P	                  ^^4      pW#3# r  )r  r  r'  r(  r:  )r  r%  r  r  s   &&  r   inc_dec_2dim_setup"TestCudaAtomics.inc_dec_2dim_setup  sV    ""2U"3yy  BR 077>FFq!L|r   c           	         VP                  4       p\        P                  ! V4      ! V4      p	WV3,          ! WV4       \        P                  P                  V\        P                  ! W8  ^ V^,           4      4       R# r2   r)  r   r*  r  ri  rj  where
r  r"   r#   r  r|  nblocksblksizer   r0  rw  s
   &&&&&&&&  r   check_inc_indexTestCudaAtomics.check_inc_index  X    xxzHHSM$'	7"#Cf5


RXXdna%JKr   c           	         VP                  4       p\        P                  ! V4      ! V4      p	WV3,          ! W!V4       \        P                  P                  V\        P                  ! W8  ^ V^,           4      4       R# r2   r  r  s
   &&&&&&&&  r   check_inc_index2 TestCudaAtomics.check_inc_index2  r	  r   c           	         VP                  4       p\        P                  ! V4      ! V4      pWV3,          ! W4       \        P                  P                  V\        P                  ! Wr8  ^ V^,           4      4       R# r2   r  	r  r"   r  r|  r  r  r   r0  rw  s	   &&&&&&&  r   	check_incTestCudaAtomics.check_inc  sV    xxzHHSM$'	7"#C0


RXXdna%JKr   c           	         V P                  \        P                  R 7      w  rpRpV P                  W#W^^ \        4       R# r$  "void(uint32[:], uint32[:], uint32)N)r  r  r   r  r   r  r  r"   r#   r|  s   &    r   test_atomic_inc_32"TestCudaAtomics.test_atomic_inc_32  9    #66RYY6G
2Sz2|Lr   c           	         V P                  \        P                  R 7      w  rpRpV P                  W#W^^ \        4       R# r$  z"void(uint64[:], uint64[:], uint64)N)r  r  r   r  r   r  s   &    r   test_atomic_inc_64"TestCudaAtomics.test_atomic_inc_64  r  r   c                ~    V P                  \        P                  4      w  rR pV P                  W!V^R\        4       R# r  NrL   )r  r  r   r  r   r  r  r"   r|  s   &   r   test_atomic_inc2_32#TestCudaAtomics.test_atomic_inc2_32  1    11"))<
)sQ~Fr   c                ~    V P                  \        P                  4      w  rR pV P                  W!V^R\        4       R# void(uint64[:,:], uint64)NrL   )r  r  r   r  r   r  s   &   r   test_atomic_inc2_64#TestCudaAtomics.test_atomic_inc2_64  r!  r   c                ~    V P                  \        P                  4      w  rR pV P                  W!V^R\        4       R# r  )r  r  r   r  r   r  s   &   r   test_atomic_inc3 TestCudaAtomics.test_atomic_inc3  1    11"))<
)sQ{Cr   c           	         V P                  \        P                  R 7      w  rpRpV P                  W#W^^ \        4       R# r  )r  r  r   r  r   r  s   &    r   test_atomic_inc_global_32)TestCudaAtomics.test_atomic_inc_global_32  <    #66RYY6G
2c
B/	1r   c           	         V P                  \        P                  R 7      w  rpRpV P                  W#W^^ \        4       R# r  )r  r  r   r  r   r  s   &    r   test_atomic_inc_global_64)TestCudaAtomics.test_atomic_inc_global_64  r.  r   c                ~    V P                  \        P                  4      w  rR pV P                  W!V^R\        4       R# r  )r  r  r   r  r   r  s   &   r   test_atomic_inc_global_2_32+TestCudaAtomics.test_atomic_inc_global_2_32  2    11"))<
)sQ7JKr   c                ~    V P                  \        P                  4      w  rR pV P                  W!V^R\        4       R# r#  )r  r  r   r  r   r  s   &   r   test_atomic_inc_global_2_64+TestCudaAtomics.test_atomic_inc_global_2_64  r5  r   c                0   VP                  4       p\        P                  ! V4      ! V4      p	WV3,          ! WV4       \        P                  P                  V\        P                  ! V^ 8H  V\        P                  ! W8  VV^,
          4      4      4       R# r2   r  r  s
   &&&&&&&&  r   check_dec_indexTestCudaAtomics.check_dec_index  u    xxzHHSM$'	7"#Cf5


RXXdai.0hht}7=7;ax/A&B 	Cr   c                0   VP                  4       p\        P                  ! V4      ! V4      p	WV3,          ! W!V4       \        P                  P                  V\        P                  ! V^ 8H  V\        P                  ! W8  VV^,
          4      4      4       R# r2   r  r  s
   &&&&&&&&  r   check_dec_index2 TestCudaAtomics.check_dec_index2  r<  r   c                .   VP                  4       p\        P                  ! V4      ! V4      pWV3,          ! W4       \        P                  P                  V\        P                  ! V^ 8H  V\        P                  ! Wr8  VV^,
          4      4      4       R# r2   r  r  s	   &&&&&&&  r   	check_decTestCudaAtomics.check_dec  ss    xxzHHSM$'	7"#C0


RXXdai.0hht}7=7;ax/A&B 	Cr   c           	         V P                  \        P                  R 7      w  rpRpV P                  W#W^^ \        4       R# r  )r  r  r   r:  r   r  s   &    r   test_atomic_dec_32"TestCudaAtomics.test_atomic_dec_32  r  r   c           	         V P                  \        P                  R 7      w  rpRpV P                  W#W^^ \        4       R# r  )r  r  r   r:  r   r  s   &    r   test_atomic_dec_64"TestCudaAtomics.test_atomic_dec_64  r  r   c                ~    V P                  \        P                  4      w  rR pV P                  W!V^R\        4       R# r  )r  r  r   rA  r   r  s   &   r   test_atomic_dec2_32#TestCudaAtomics.test_atomic_dec2_32  r!  r   c                ~    V P                  \        P                  4      w  rR pV P                  W!V^R\        4       R# r#  )r  r  r   rA  r   r  s   &   r   test_atomic_dec2_64#TestCudaAtomics.test_atomic_dec2_64  r!  r   c                ~    V P                  \        P                  4      w  rR pV P                  W!V^R\        4       R# r  )r  r  r   rA  r   r  s   &   r   test_atomic_dec3_new$TestCudaAtomics.test_atomic_dec3_new  r*  r   c           	         V P                  \        P                  R 7      w  rpRpV P                  W#W^^ \        4       R# r  )r  r  r   r>  r   r  s   &    r   test_atomic_dec_global_32)TestCudaAtomics.test_atomic_dec_global_32!  r.  r   c           	         V P                  \        P                  R 7      w  rpRpV P                  W#W^^ \        4       R# r  )r  r  r   r>  r   r  s   &    r   test_atomic_dec_global_64)TestCudaAtomics.test_atomic_dec_global_64'  r.  r   c                ~    V P                  \        P                  4      w  rR pV P                  W!V^R\        4       R# r  )r  r  r   rA  r   r  s   &   r   test_atomic_dec_global2_32*TestCudaAtomics.test_atomic_dec_global2_32-  r5  r   c                ~    V P                  \        P                  4      w  rR pV P                  W!V^R\        4       R# r#  )r  r  r   rA  r   r  s   &   r   test_atomic_dec_global2_64*TestCudaAtomics.test_atomic_dec_global2_642  r5  r   c                   \         P                  P                  ^2^d\         P                  R7      p\         P                  P                  ^ ^ ^ R7      P	                  \         P                  4      p\         P
                  ! ^ \         P                  R7      p\        P                  ! R4      ! \        4      pVR,          ! W#V4       \         P                  P                  W!4       R# )2   r$  r#  r  Nr&  )r  r  r'  r   r(  r  r   r*  r   ri  rj  )r  r  r"   r#   rw  s   &    r   test_atomic_exch TestCudaAtomics.test_atomic_exch7  s    YY&&r3bii&@
ii2B/66ryyAii")),HHAB;O	%:.


0r   c                   \         P                  P                  ^2^d\         P                  R7      p\         P                  P                  ^ ^ ^ R7      P	                  \         P                  4      P                  ^^4      p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  P                  W!4       R# )r_  r$  r#  r  Nr9  )r  r  r'  r   r(  r:  r   r*  r   ri  rj  r  r  r"   rw  s   &   r   test_atomic_exch2!TestCudaAtomics.test_atomic_exch2A      YY&&r3bii&@
ii2B/66ryyAII!QOHH89,G	)S-


0r   c                   \         P                  P                  ^2^d\         P                  R7      p\         P                  P                  ^ ^ ^ R7      P	                  \         P                  4      P                  ^^4      p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  P                  W!4       R# )r_  r$  r#  r$  Nr9  )r  r  r'  r   r(  r:  r   r*  r   ri  rj  rc  s   &   r   test_atomic_exch3!TestCudaAtomics.test_atomic_exch3I  rf  r   c                   \         P                  P                  ^2^d\         P                  R7      p\         P                  ! ^ \         P                  R7      p\         P                  P                  ^ ^ ^ \         P                  R7      pRp\
        P                  ! V4      ! \        4      pVR,          ! W#V4       \         P                  P                  W14       R# )r_  r$  rf  r  Nr&  )
r  r  r'  r   r  r   r*  r   ri  rj  )r  r  r#   r"   r|  rw  s   &     r   test_atomic_exch_global'TestCudaAtomics.test_atomic_exch_globalQ  s    YY&&r3bii&@
ii")),ii2Bbii@2HHSM"45	%:.


0r   c                f   \         P                  P                  W#RR7      P                  V4      p\         P                  ! ^VP
                  R7      p\        P                  ! \        4      pVR,          ! WT4       \         P                  ! V4      p\         P                  P                  WW4       R# )ro   r#  r$  Nro   ro   )r  r  r'  r(  r+  r%  r   r*  
atomic_maxmaxri  rj  r  r%  lohivalsr  rw  r3  s   &&&&    r   check_atomic_max TestCudaAtomics.check_atomic_max[  st    yy  h 7>>uEhhq

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


*r   c                J    V P                  \        P                  RR R7       R#   r%  rr  rs  N )ru  r  r  r  s   &r   test_atomic_max_int32%TestCudaAtomics.test_atomic_max_int32c      BHHEBr   c                J    V P                  \        P                  ^ RR7       R# r   ry  rz  N)ru  r  r   r|  s   &r   test_atomic_max_uint32&TestCudaAtomics.test_atomic_max_uint32f      BII!>r   c                J    V P                  \        P                  RR R7       R# rx  )ru  r  rh  r|  s   &r   test_atomic_max_int64%TestCudaAtomics.test_atomic_max_int64i  r  r   c                J    V P                  \        P                  ^ RR7       R# r  )ru  r  r   r|  s   &r   test_atomic_max_uint64&TestCudaAtomics.test_atomic_max_uint64l  r  r   c                J    V P                  \        P                  RR R7       R# rx  )ru  r  r   r|  s   &r   test_atomic_max_float32'TestCudaAtomics.test_atomic_max_float32o      BJJ6eDr   c                J    V P                  \        P                  RR R7       R# rx  )ru  r  r   r|  s   &r   test_atomic_max_double&TestCudaAtomics.test_atomic_max_doubler  r  r   c                   \         P                  P                  ^ RRR7      P                  \         P                  4      p\         P
                  ! ^\         P                  4      p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  ! V4      p\         P                  P                  W$4       R# r   ry  r#  void(float64[:], float64[:,:])Nrn  )r  r  r'  r(  r   r+  r   r*  !atomic_max_double_normalizedindexrp  ri  rj  r  rt  r  rw  r3  s   &    r   &test_atomic_max_double_normalizedindex6TestCudaAtomics.test_atomic_max_double_normalizedindexu  s    yy  E 9@@Lhhq"**%HH=>-/	&#$vvd|


*r   c                   \         P                  P                  ^ ^^ R7      P                  \         P                  4      p\         P
                  ! ^\         P                  4      p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  ! V4      p\         P                  P                  W$4       R# r   r#  void(float64[:], float64[:])Nr&  )r  r  r'  r(  r   r+  r   r*  atomic_max_double_oneindexrp  ri  rj  r  s   &    r   test_atomic_max_double_oneindex/TestCudaAtomics.test_atomic_max_double_oneindex  s    yy  Cb 188Dhhq"**%HH;<&(	%#vvd|


*r   c                h   \         P                  P                  W#RR7      P                  V4      p\         P                  ! R.VP
                  R7      p\        P                  ! \        4      pVR,          ! WT4       \         P                  ! V4      p\         P                  P                  WW4       R# )ro   r#  ry  r$  Nrn  )r  r  r'  r(  r    r%  r   r*  
atomic_minminri  rj  rq  s   &&&&    r   check_atomic_min TestCudaAtomics.check_atomic_min  sv    yy  h 7>>uEhhwdjj1HHZ(	&#$vvd|


*r   c                J    V P                  \        P                  RR R7       R# rx  )r  r  r  r|  s   &r   test_atomic_min_int32%TestCudaAtomics.test_atomic_min_int32  r  r   c                J    V P                  \        P                  ^ RR7       R# r  )r  r  r   r|  s   &r   test_atomic_min_uint32&TestCudaAtomics.test_atomic_min_uint32  r  r   c                J    V P                  \        P                  RR R7       R# rx  )r  r  rh  r|  s   &r   test_atomic_min_int64%TestCudaAtomics.test_atomic_min_int64  r  r   c                J    V P                  \        P                  ^ RR7       R# r  )r  r  r   r|  s   &r   test_atomic_min_uint64&TestCudaAtomics.test_atomic_min_uint64  r  r   c                J    V P                  \        P                  RR R7       R# rx  )r  r  r   r|  s   &r   test_atomic_min_float%TestCudaAtomics.test_atomic_min_float  r  r   c                J    V P                  \        P                  RR R7       R# rx  )r  r  r   r|  s   &r   test_atomic_min_double&TestCudaAtomics.test_atomic_min_double  r  r   c                   \         P                  P                  ^ RRR7      P                  \         P                  4      p\         P
                  ! ^\         P                  4      R,          p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  ! V4      p\         P                  P                  W$4       R# r  )r  r  r'  r(  r   onesr   r*  !atomic_min_double_normalizedindexr  ri  rj  r  s   &    r   &test_atomic_min_double_normalizedindex6TestCudaAtomics.test_atomic_min_double_normalizedindex  s    yy  E 9@@Lgga$u,HH=>-/	&#$vvd|


*r   c                   \         P                  P                  ^ ^^ R7      P                  \         P                  4      p\         P
                  ! ^\         P                  4      ^,          p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  ! V4      p\         P                  P                  W$4       R# r  )r  r  r'  r(  r   r  r   r*  atomic_min_double_oneindexr  ri  rj  r  s   &    r   test_atomic_min_double_oneindex/TestCudaAtomics.test_atomic_min_double_oneindex  s    yy  Cb 188Dgga$s*HH;<&(	%#vvd|


*r   c                   \         P                  ! R 4      ! V4      p\        P                  P	                  ^ ^RR7      P                  \        P                  4      p\        P                  ! ^\        P                  4      \        P                  ,           pVR,          ! WC4       \        P                  P                  V\        P                  .4       R# r  r#  Nr@   r@   )r   r*  r  r  r'  r(  r   r+  nanri  rj  )r  r   rw  rt  r  s   &&   r    _test_atomic_minmax_nan_location0TestCudaAtomics._test_atomic_minmax_nan_location  s    HH=>tD	yy  Ce 4;;BJJGhhq"**%.$"


bffX.r   c                   \         P                  ! R 4      ! V4      p\        P                  P	                  ^ ^^R7      P                  \        P                  4      pVP                  4       p\        P                  ! R\        P                  4      \        P                  ,           pVR,          ! W54       \        P                  P                  W44       R# r  )r   r*  r  r  r'  r(  r   r)  r+  r  ri  rj  )r  r   rw  r  r3  rt  s   &&    r   _test_atomic_minmax_nan_val+TestCudaAtomics._test_atomic_minmax_nan_val  s    HH=>tD	ii3Q/66rzzBxxzxx

+bff4$"


*r   c                0    V P                  \        4       R # r   )r  r  r|  s   &r   test_atomic_min_nan_location,TestCudaAtomics.test_atomic_min_nan_location      --j9r   c                0    V P                  \        4       R # r   )r  ro  r|  s   &r   test_atomic_max_nan_location,TestCudaAtomics.test_atomic_max_nan_location  r  r   c                0    V P                  \        4       R # r   )r  r  r|  s   &r   test_atomic_min_nan_val'TestCudaAtomics.test_atomic_min_nan_val      ((4r   c                0    V P                  \        4       R # r   )r  ro  r|  s   &r   test_atomic_max_nan_val'TestCudaAtomics.test_atomic_max_nan_val  r  r   c                   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      p\         P
                  ! ^\         P                  4      pRp\        P                  ! V4      ! \        4      pVR,          ! W!4       \         P                  ! V4      p\         P                  P                  W%4       R# r  )r  r  r'  r(  r   r+  r   r*  atomic_max_double_sharedrp  ri  rj  r  rt  r  r|  rw  r3  s   &     r   test_atomic_max_double_shared-TestCudaAtomics.test_atomic_max_double_shared  s    yy  BR 077

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


*r   c                   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      p\         P
                  ! ^\         P                  4      ^ ,          pRp\        P                  ! V4      ! \        4      pVR,          ! W!4       \         P                  ! V4      p\         P                  P                  W%4       R# r  )r  r  r'  r(  r   r  r   r*  atomic_min_double_sharedr  ri  rj  r  s   &     r   test_atomic_min_double_shared-TestCudaAtomics.test_atomic_min_double_shared  s    yy  BR 077

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


*r   c                   V.V^,          ,          V.V^,          ,          ,           p\         P                  P                  V4       \         P                  ! WtR7      pV^8X  d   VP	                  R4      p\         P
                  ! V4      p\         P                  P                  ^^
VP                  R7      P                  VP                  4      p	Wr8H  p
Ws8H  p\         P
                  ! V4      pW,          W&   W<V&   VP                  4       p\        P                  ! V4      pV^8X  d   VR,          ! WxW4       MVR,          ! WxW4       \         P                  P                  W4       \         P                  P                  W4       R# )r  r$  r#  N)
   r  r  )r  r  )r  r  shuffleasarrayr:  
zeros_liker'  r:   r(  r%  r)  r   r*  ri  assert_array_equal)r  nfillunfillr%  cas_funcndimr  outr"   	fill_maskunfill_mask
expect_res
expect_outrw  s   &&&&&&&        r   	check_casTestCudaAtomics.check_cas  s$   fQ6(a1f"55
		#jj*19++h'CmmC ii2CII6==ciiHK	m]]3'
 #
"(;XXZ
HHX&	19fc2()#C>


%%j6


%%j6r   c                V    V P                  ^dRR\        P                  \        R7       R# d   r  r  r  r%  r  Nr  )r  r  r  r  r|  s   &r   test_atomic_compare_and_swap,TestCudaAtomics.test_atomic_compare_and_swap  "    3r 7 	 	9r   c                V    V P                  ^dRR\        P                  \        R7       R# r  r  Nr  )r  r  rh  r  r|  s   &r   test_atomic_compare_and_swap2-TestCudaAtomics.test_atomic_compare_and_swap2  r  r   c                   \         P                  P                  ^2R\         P                  R7      p\         P                  P                  ^^\         P                  R7      pV P	                  ^dW\         P                  \
        R7       R# r_  r  r$  r  N)r  r  r'  r   r  r  r  rfillrunfills   &  r   test_atomic_compare_and_swap3-TestCudaAtomics.test_atomic_compare_and_swap3  ^    		!!"c!;))##Ar#;5		 7 	 	9r   c                   \         P                  P                  ^2R\         P                  R7      p\         P                  P                  ^^\         P                  R7      pV P	                  ^dW\         P                  \
        R7       R# r  )r  r  r'  r   r  r  r  s   &  r   test_atomic_compare_and_swap4-TestCudaAtomics.test_atomic_compare_and_swap4  r  r   c                V    V P                  ^dRR\        P                  \        R7       R# r  )r  r  r  r  r|  s   &r   test_atomic_cas_1dim$TestCudaAtomics.test_atomic_cas_1dim%  "    3r / 	 	1r   c           	     X    V P                  ^dRR\        P                  \        ^R7       R# )r  r  r  r  r%  r  r  Nr  r  )r  r  r  r  r|  s   &r   test_atomic_cas_2dim$TestCudaAtomics.test_atomic_cas_2dim)  $    3r /a 	 	9r   c                V    V P                  ^dRR\        P                  \        R7       R# r  )r  r  rh  r  r|  s   &r   test_atomic_cas2_1dim%TestCudaAtomics.test_atomic_cas2_1dim-  r  r   c           	     X    V P                  ^dRR\        P                  \        ^R7       R# )r  r  Nr  r  )r  r  rh  r  r|  s   &r   test_atomic_cas2_2dim%TestCudaAtomics.test_atomic_cas2_2dim1  r  r   c                   \         P                  P                  ^2R\         P                  R7      p\         P                  P                  ^^\         P                  R7      pV P	                  ^dW\         P                  \
        R7       R# r  )r  r  r'  r   r  r  r  s   &  r   test_atomic_cas3_1dim%TestCudaAtomics.test_atomic_cas3_1dim5  ^    		!!"c!;))##Ar#;5		 / 	 	1r   c           	        \         P                  P                  ^2R\         P                  R7      p\         P                  P                  ^^\         P                  R7      pV P	                  ^dW\         P                  \
        ^R7       R# r_  r  r$  r  N)r  r  r'  r   r  r  r  s   &  r   test_atomic_cas3_2dim%TestCudaAtomics.test_atomic_cas3_2dim;  `    		!!"c!;))##Ar#;5		 /a 	 	9r   c                   \         P                  P                  ^2R\         P                  R7      p\         P                  P                  ^^\         P                  R7      pV P	                  ^dW\         P                  \
        R7       R# r  )r  r  r'  r   r  r  r  s   &  r   test_atomic_cas4_1dim%TestCudaAtomics.test_atomic_cas4_1dimA  r  r   c           	        \         P                  P                  ^2R\         P                  R7      p\         P                  P                  ^^\         P                  R7      pV P	                  ^dW\         P                  \
        ^R7       R# r  )r  r  r'  r   r  r  r  s   &  r   test_atomic_cas4_2dim%TestCudaAtomics.test_atomic_cas4_2dimG  r"  r   c                >   \         P                  ! ^\         P                  R7      pW#^ &   VR,          ! V4       \         P                  ! V4      '       d/   V P	                  \         P                  ! V^,          4      4       R# V P                  V^,          V4       R# r  r$  Nr  )r  r+  r   isnanr-  assertEqualr  r_  initialr   s   &&& r   _test_atomic_returns_old(TestCudaAtomics._test_atomic_returns_oldR  sc    HHQbjj)!tQ88GOOBHHQqTN+QqT7+r   c                V    \         P                  R  4       pV P                  V^
4       R# )c                 N    \         P                  P                  V ^ ^4      V ^&   R# r2   )r   rB   rC   r   s   &r   r_  ;TestCudaAtomics.test_atomic_add_returns_old.<locals>.kernel\      ;;??1a+AaDr   Nr   r*  r/  r  r_  s   & r   test_atomic_add_returns_old+TestCudaAtomics.test_atomic_add_returns_old[  *    		, 
	, 	%%fb1r   c                V    \         P                  R  4       pV P                  V^
4       R# )c                 N    \         P                  P                  V ^ ^4      V ^&   R# r2   r   rB   rp  r3  s   &r   r_  BTestCudaAtomics.test_atomic_max_returns_no_replace.<locals>.kernelc  r5  r   Nr6  r7  s   & r   "test_atomic_max_returns_no_replace2TestCudaAtomics.test_atomic_max_returns_no_replaceb  r:  r   c                V    \         P                  R  4       pV P                  V^4       R# )c                 N    \         P                  P                  V ^ ^
4      V ^&   R# r2   r=  r3  s   &r   r_  CTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernelj      ;;??1a,AaDr   Nr6  r7  s   & r   #test_atomic_max_returns_old_replace3TestCudaAtomics.test_atomic_max_returns_old_replacei  s*    		- 
	- 	%%fa0r   c                r    \         P                  R  4       pV P                  V\        P                  4       R# )c                 N    \         P                  P                  V ^ ^4      V ^&   R# r2   r=  r3  s   &r   r_  HTestCudaAtomics.test_atomic_max_returns_old_nan_in_array.<locals>.kernelq  r5  r   Nr   r*  r/  r  r  r7  s   & r   (test_atomic_max_returns_old_nan_in_array8TestCudaAtomics.test_atomic_max_returns_old_nan_in_arrayp  s.    		, 
	, 	%%fbff5r   c                V    \         P                  R  4       pV P                  V^
4       R# )c                 j    \         P                  P                  V ^ \        P                  4      V ^&   R# r2   )r   rB   rp  r  r  r3  s   &r   r_  CTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernelx       ;;??1a0AaDr   Nr6  r7  s   & r   #test_atomic_max_returns_old_nan_val3TestCudaAtomics.test_atomic_max_returns_old_nan_valw  *    		1 
	1 	%%fb1r   c                V    \         P                  R  4       pV P                  V^
4       R# )c                 N    \         P                  P                  V ^ ^4      V ^&   R# r2   r   rB   r  r3  s   &r   r_  FTestCudaAtomics.test_atomic_min_returns_old_no_replace.<locals>.kernel  rD  r   Nr6  r7  s   & r   &test_atomic_min_returns_old_no_replace6TestCudaAtomics.test_atomic_min_returns_old_no_replace~  *    		- 
	- 	%%fb1r   c                V    \         P                  R  4       pV P                  V^4       R# )c                 N    \         P                  P                  V ^ ^
4      V ^&   R# r2   rV  r3  s   &r   r_  CTestCudaAtomics.test_atomic_min_returns_old_replace.<locals>.kernel  rD  r   Nr6  r7  s   & r   #test_atomic_min_returns_old_replace3TestCudaAtomics.test_atomic_min_returns_old_replace  rZ  r   c                r    \         P                  R  4       pV P                  V\        P                  4       R# )c                 N    \         P                  P                  V ^ ^4      V ^&   R# r2   rV  r3  s   &r   r_  HTestCudaAtomics.test_atomic_min_returns_old_nan_in_array.<locals>.kernel  rD  r   NrJ  r7  s   & r   (test_atomic_min_returns_old_nan_in_array8TestCudaAtomics.test_atomic_min_returns_old_nan_in_array  s.    		- 
	- 	%%fbff5r   c                V    \         P                  R  4       pV P                  V^4       R# )c                 j    \         P                  P                  V ^ \        P                  4      V ^&   R# r2   )r   rB   r  r  r  r3  s   &r   r_  CTestCudaAtomics.test_atomic_min_returns_old_nan_val.<locals>.kernel  rP  r   Nr6  r7  s   & r   #test_atomic_min_returns_old_nan_val3TestCudaAtomics.test_atomic_min_returns_old_nan_val  rS  r   c                n   \         P                  P                  W#RR7      P                  V4      pWER&   \         P                  ! ^VP
                  R7      p\        P                  ! \        4      pVR,          ! We4       \         P                  ! V4      p\         P                  P                  Wh4       R# )ro   r#  r@   Nr  r$  Nrn  )r  r  r'  r(  r+  r%  r   r*  atomic_nanmaxnanmaxri  rj  	r  r%  rr  rs  init_valrt  r  rw  r3  s	   &&&&&    r   check_atomic_nanmax#TestCudaAtomics.check_atomic_nanmax  s{    yy  h 7>>uET
hhq

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


*r   c                L    V P                  \        P                  RR ^ R7       R# ry  r%  rr  rs  ro  Nr{  )rp  r  r  r|  s   &r   test_atomic_nanmax_int32(TestCudaAtomics.test_atomic_nanmax_int32  "      rxxFu*+ 	! 	-r   c                L    V P                  \        P                  ^ R^ R7       R# r   ry  rt  N)rp  r  r   r|  s   &r   test_atomic_nanmax_uint32)TestCudaAtomics.test_atomic_nanmax_uint32  "      ryyQ5*+ 	! 	-r   c                L    V P                  \        P                  RR ^ R7       R# rs  )rp  r  rh  r|  s   &r   test_atomic_nanmax_int64(TestCudaAtomics.test_atomic_nanmax_int64  rw  r   c                L    V P                  \        P                  ^ R^ R7       R# ry  )rp  r  r   r|  s   &r   test_atomic_nanmax_uint64)TestCudaAtomics.test_atomic_nanmax_uint64  r|  r   c                h    V P                  \        P                  RR \        P                  R7       R# rs  )rp  r  r   r  r|  s   &r   test_atomic_nanmax_float32*TestCudaAtomics.test_atomic_nanmax_float32  &      rzzf*,&& 	! 	2r   c                h    V P                  \        P                  RR \        P                  R7       R# rs  )rp  r  r   r  r|  s   &r   test_atomic_nanmax_double)TestCudaAtomics.test_atomic_nanmax_double  r  r   c                   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      p\         P
                  VR&   \         P                  ! ^ .VP                  R7      pRp\        P                  ! V4      ! \        4      pVR,          ! W!4       \         P                  ! V4      p\         P                  P                  W%4       R# r   r#  rk  r$  r  Nr&  )r  r  r'  r(  r   r  r    r%  r   r*  atomic_nanmax_double_sharedrm  ri  rj  r  s   &     r    test_atomic_nanmax_double_shared0TestCudaAtomics.test_atomic_nanmax_double_shared  s    yy  BR 077

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


*r   c                   \         P                  P                  ^ ^^ R7      P                  \         P                  4      p\         P
                  VR&   \         P                  ! ^\         P                  4      p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  ! V4      p\         P                  P                  W$4       R# r   r#  rk  r  Nr&  )r  r  r'  r(  r   r  r+  r   r*  r  rm  ri  rj  r  s   &    r   "test_atomic_nanmax_double_oneindex2TestCudaAtomics.test_atomic_nanmax_double_oneindex  s    yy  Cb 188DVVT
hhq"**%HH;<&(	%#yy


*r   c                p   \         P                  P                  W#RR7      P                  V4      pWER&   \         P                  ! R.VP
                  R7      p\        P                  ! \        4      pVR,          ! We4       \         P                  ! V4      p\         P                  P                  Wh4       R# )ro   r#  rk  ry  r$  Nrn  )r  r  r'  r(  r    r%  r   r*  atomic_nanminnanminri  rj  rn  s	   &&&&&    r   check_atomic_nanmin#TestCudaAtomics.check_atomic_nanmin  s}    yy  h 7>>uET
hhwdjj1HH]+	&#$yy


*r   c                L    V P                  \        P                  RR ^ R7       R# rs  )r  r  r  r|  s   &r   test_atomic_nanmin_int32(TestCudaAtomics.test_atomic_nanmin_int32  rw  r   c                L    V P                  \        P                  ^ R^ R7       R# ry  )r  r  r   r|  s   &r   test_atomic_nanmin_uint32)TestCudaAtomics.test_atomic_nanmin_uint32  r|  r   c                L    V P                  \        P                  RR ^ R7       R# rs  )r  r  rh  r|  s   &r   test_atomic_nanmin_int64(TestCudaAtomics.test_atomic_nanmin_int64  rw  r   c                L    V P                  \        P                  ^ R^ R7       R# ry  )r  r  r   r|  s   &r   test_atomic_nanmin_uint64)TestCudaAtomics.test_atomic_nanmin_uint64  r|  r   c                h    V P                  \        P                  RR \        P                  R7       R# rs  )r  r  r   r  r|  s   &r   test_atomic_nanmin_float(TestCudaAtomics.test_atomic_nanmin_float  r  r   c                h    V P                  \        P                  RR \        P                  R7       R# rs  )r  r  r   r  r|  s   &r   test_atomic_nanmin_double)TestCudaAtomics.test_atomic_nanmin_double  r  r   c                   \         P                  P                  ^ ^ ^ R7      P                  \         P                  4      p\         P
                  VR&   \         P                  ! ^ .VP                  R7      pRp\        P                  ! V4      ! \        4      pVR,          ! W!4       \         P                  ! V4      p\         P                  P                  W%4       R# r  )r  r  r'  r(  r   r  r    r%  r   r*  atomic_nanmin_double_sharedr  ri  rj  r  s   &     r    test_atomic_nanmin_double_shared0TestCudaAtomics.test_atomic_nanmin_double_shared  s    yy  BR 077

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


*r   c                   \         P                  P                  ^ ^^ R7      P                  \         P                  4      p\         P
                  VR&   \         P                  ! ^.\         P                  4      p\        P                  ! R4      ! \        4      pVR,          ! W!4       \         P                  ! V4      p\         P                  P                  W$4       R# r  )r  r  r'  r(  r   r  r    r   r*  r  r  ri  rj  r  s   &    r   "test_atomic_nanmin_double_oneindex2TestCudaAtomics.test_atomic_nanmin_double_oneindex  s    yy  Cb 188DVVT
hhubjj)HH;<&(	%#yy


*r   c                   \         P                  ! ^\         P                  R7      pW#^ &   \         P                  V^&   VR,          ! V4       \         P                  ! V4      '       d[   V P                  \         P                  ! V^ ,          4      4       V P                  \         P                  ! V^,          4      4       R# V P                  V^,          V4       R# r*  )r  r+  r   r  r+  assertFalser-  r,  r-  s   &&& r   _test_atomic_nan_returns_old,TestCudaAtomics._test_atomic_nan_returns_old  s    HHQbjj)!vv!tQ88GRXXad^,OOBHHQqTN+QqT7+r   c                V    \         P                  R  4       pV P                  V^
4       R# )c                 N    \         P                  P                  V ^ ^4      V ^&   R# r2   r   rB   rm  r3  s   &r   r_  ITestCudaAtomics.test_atomic_nanmax_returns_old_no_replace.<locals>.kernel      ;;%%aA.AaDr   Nr   r*  r  r7  s   & r   )test_atomic_nanmax_returns_old_no_replace9TestCudaAtomics.test_atomic_nanmax_returns_old_no_replace  s*    		/ 
	/ 	))&"5r   c                V    \         P                  R  4       pV P                  V^4       R# )c                 N    \         P                  P                  V ^ ^
4      V ^&   R# r2   r  r3  s   &r   r_  FTestCudaAtomics.test_atomic_nanmax_returns_old_replace.<locals>.kernel"      ;;%%aB/AaDr   Nr  r7  s   & r   &test_atomic_nanmax_returns_old_replace6TestCudaAtomics.test_atomic_nanmax_returns_old_replace!  s*    		0 
	0 	))&!4r   c                r    \         P                  R  4       pV P                  V\        P                  4       R# )c                 N    \         P                  P                  V ^ ^4      V ^&   R# r2   r  r3  s   &r   r_  KTestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array.<locals>.kernel)  r  r   Nr   r*  r  r  r  r7  s   & r   +test_atomic_nanmax_returns_old_nan_in_array;TestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array(  s.    		/ 
	/ 	))&"&&9r   c                V    \         P                  R  4       pV P                  V^
4       R# )c                 j    \         P                  P                  V ^ \        P                  4      V ^&   R# r2   )r   rB   rm  r  r  r3  s   &r   r_  FTestCudaAtomics.test_atomic_nanmax_returns_old_nan_val.<locals>.kernel0  "    ;;%%aBFF3AaDr   Nr  r7  s   & r   &test_atomic_nanmax_returns_old_nan_val6TestCudaAtomics.test_atomic_nanmax_returns_old_nan_val/  *    		4 
	4 	))&"5r   c                V    \         P                  R  4       pV P                  V^
4       R# )c                 N    \         P                  P                  V ^ ^4      V ^&   R# r2   r   rB   r  r3  s   &r   r_  ITestCudaAtomics.test_atomic_nanmin_returns_old_no_replace.<locals>.kernel7  r  r   Nr  r7  s   & r   )test_atomic_nanmin_returns_old_no_replace9TestCudaAtomics.test_atomic_nanmin_returns_old_no_replace6  *    		0 
	0 	))&"5r   c                V    \         P                  R  4       pV P                  V^4       R# )c                 N    \         P                  P                  V ^ ^
4      V ^&   R# r2   r  r3  s   &r   r_  FTestCudaAtomics.test_atomic_nanmin_returns_old_replace.<locals>.kernel>  r  r   Nr  r7  s   & r   &test_atomic_nanmin_returns_old_replace6TestCudaAtomics.test_atomic_nanmin_returns_old_replace=  r  r   c                r    \         P                  R  4       pV P                  V\        P                  4       R# )c                 N    \         P                  P                  V ^ ^4      V ^&   R# r2   r  r3  s   &r   r_  KTestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_array.<locals>.kernelE  r  r   Nr  r7  s   & r   +test_atomic_nanmin_returns_old_nan_in_array;TestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_arrayD  s.    		0 
	0 	))&"&&9r   c                V    \         P                  R  4       pV P                  V^4       R# )c                 j    \         P                  P                  V ^ \        P                  4      V ^&   R# r2   )r   rB   r  r  r  r3  s   &r   r_  FTestCudaAtomics.test_atomic_nanmin_returns_old_nan_val.<locals>.kernelL  r  r   Nr  r7  s   & r   &test_atomic_nanmin_returns_old_nan_val6TestCudaAtomics.test_atomic_nanmin_returns_old_nan_valK  r  r   r   )T)r@   )__name__
__module____qualname____firstlineno__r  r5  r=  rB  rI  rN  rQ  rb  rm  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,  r0  r3  r7  r:  r>  rA  rD  rG  rJ  rM  rP  rS  rV  rY  r\  r`  rd  rh  rk  ru  r}  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r	  r  r  r  r  r  r   r$  r'  r/  r8  r?  rE  rK  rQ  rX  r^  rc  rh  rp  ru  rz  r~  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  __static_attributes____classdictcell____classcell__)r  __classdict__s   @@r   r  r    s    2$612$61322(7 4G*G"B
-11
-11
+//+//-::+8-::+8-::+8%
LLLM
M
G
G
D
11L
L
CCCM
M
G
G
D
11L
L
1111+C?C?EE+++C?C?EE++*/+::55++76999919191919,221622262+----22	+	++----22	+	+	,65:666:6 6r   r  __main__)qnumpyr  textwrapr   numbar   r   r   r   r   numba.cuda.testingr	   r
   r   
numba.corer   r*  r   r   r   r.   r0   r8   r;   r=   rE   rI   rP   rS   rU   r\   r`   rd   rh   rk   rr   ru   rx   rz   r|   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  ro  r  r  r  r  r  r  r  rl  $atomic_nanmax_double_normalizedindexatomic_nanmax_double_oneindexr  r  $atomic_nanmin_double_normalizedindexatomic_nanmin_double_oneindexr  r  r  r  r  r  mainr   r   r   <module>r     s+     8 8 D D          	 	      K
J
H
G
M
O
N
H
G
M
IHPO%
M
L
H
G
M
K
H
M
O
H
M
M
H
M
I%
M
L
I
N
J7
K
H
M
I7
K
H
M
I7
B
D
H
H
M
I7
B
D
H
H
M
I7
C
I
I
J!FJ 66GH.0J56GH.0J 12=4 ; 12=4 ;OAA{6l {6|# zMMO r   