
    shQ                         S SK rS SKJr  S SKJrJrJrJrJ	r	  S SK
JrJrJr  S SKJr  \R                   " SS9S 5       r\R                   " SS9S	 5       r\R                   " SS9S
 5       r\R                   " SS9S 5       r\R                   " SS9S 5       r\R                   " SS9S 5       r\R                   " SS9S 5       r\R                   " SS9S 5       rS rS rS rS rS rS rS rS r S r!S r"S r#S r$S r%S r&S r'S r(S  r)S! r*S" r+S# r,S$ r-S% r.S& r/S' r0S( r1S) r2S* r3S+ r4S, r5S- r6S. r7S/ r8S0 r9S1 r:S2 r;S3 r<S4 r=S5 r>S6 r?S7 r@S8 rAS9 rBS: rCS; rDS< rES= rFS> rGS? rHS@ rISA rJSB rKSC rLSD rMSE rNSF rOSG rPSH rQSI rRSJ rSSK rTSL rUSM rVSN rWSO rXSP rYSQ rZ\Z" SR5      u  r[r\r]r^\Z" SS5      u  r_r`rarb\Z" ST5      u  rcrdrerf\Z" SU5      u  rgrhrirjSV rkSW rlSX rm " SY SZ\5      rn\oS[:X  a  \R                  " 5         gg)\    N)dedent)cudauint32uint64float32float64)unittestCUDATestCasecc_X_or_above)configT)devicec                     [        U 5      $ N)r   nums    x/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_atomics.pyatomic_cast_to_uint64r   	   s    #;    c                     [        U 5      $ r   )intr   s    r   atomic_cast_to_intr      s    s8Or   c                     U $ r    r   s    r   atomic_cast_noner      s    Jr   c	                     [         R                  R                  n	[         R                  R	                  XC5      n
XzU	'   [         R
                  " 5         U" X   U-  5      nU(       a  X-  nU" XU5        [         R
                  " 5         X   X	'   g 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                    [         R                  R                  n[         R                  R	                  XC5      nX   X'   [         R
                  " 5         U" X   U-  5      n	U" XU5        [         R
                  " 5         X   X'   g 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                    [         R                  R                  n[         R                  R                  n[         R                  R                  X25      n	XU4   XU4'   [         R                  " 5         Xu" U5      4n
U(       a  U
S   US   -  U
S   US   -  4n
U" XU5        [         R                  " 5         XU4   XU4'   g Nr      )r   r   r   yr   r    r!   )r"   r$   r%   	ary_shaper'   y_cast_funcr*   txtyr,   r-   s              r   atomic_binary_2dim_sharedr9   5   s     
		B			B			9	0BVB2vJ{2
C1v	!$c!fy|&;<rV*CBKr   c                     [         R                  R                  n[         R                  R                  nXS" U5      4nU(       a*  US   U R                  S   -  US   U R                  S   -  4nU" XU5        g r2   )r   r   r   r4   shape)r"   r$   r'   r6   r*   r7   r8   r-   s           r   atomic_binary_2dim_globalr<   E   sg    			B			B{2
C1v		!$c!fsyy|&;<sr   c                     [         R                  R                  n[        X   U-  5      nU(       a  Xr-  nU" XU5        g 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                 h    [        X S[        S[        R                  R                  [
        SS5	        g Nr3       r   Fr.   r   r   atomicaddr   r"   s    r   
atomic_addrF   Y   %    c62"kkoo/?EKr   c                 h    [        X S[        S[        R                  R                  [
        SS5	        g )Nr3   rA   r   TrB   rE   s    r   atomic_add_wraprI   ^   s%    c62"kkoo/?DJr   c           	      f    [        U S[        S[        R                  R                  [
        S5        g Nr3         Fr9   r   r   rC   rD   r   rE   s    r   atomic_add2rP   c   #    c1ff"kkoo/?Hr   c           	      f    [        U S[        S[        R                  R                  [
        S5        g )Nr3   rL   TrO   rE   s    r   atomic_add2_wraprS   h   s#    c1ff"kkoo/?Gr   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rK   )r9   r   r   rC   rD   r   rE   s    r   atomic_add3rU   m   #    c1ff"kkoo/DeMr   c                 h    [        X S[        S[        R                  R                  [
        SS5	        g N      ?rA           Fr.   r   r   rC   rD   r   rE   s    r   atomic_add_floatr\   r   %    cWb"kkoo/A3Or   c                 h    [        X S[        S[        R                  R                  [
        SS5	        g NrY   rA   rZ   Tr[   rE   s    r   atomic_add_float_wrapr`   w   s%    cWb"kkoo/A3Nr   c           	      f    [        U S[        S[        R                  R                  [
        S5        g NrY   rL   Fr9   r   r   rC   rD   r   rE   s    r   atomic_add_float_2rd   |   #    c3"kkoo/?Hr   c           	      f    [        U S[        S[        R                  R                  [
        S5        g NrY   rL   Trc   rE   s    r   atomic_add_float_2_wraprh      #    c3"kkoo/?Gr   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rb   )r9   r   r   rC   rD   r   rE   s    r   atomic_add_float_3rk      #    c3"kkoo/DeMr   c                 R    [        XSS[        R                  R                  S5        g NrA   rY   Fr>   r   rC   rD   r#   r"   s     r   atomic_add_double_globalrq          cC%Hr   c                 R    [        XSS[        R                  R                  S5        g )NrA   rY   Tro   rp   s     r   atomic_add_double_global_wraprt      s    cC$Gr   c                 Z    [        U S[        R                  R                  [        S5        g Nr3   Fr<   r   rC   rD   r   rE   s    r   atomic_add_double_global_2rx      s    c1dkkoo7GOr   c                 Z    [        U S[        R                  R                  [        S5        g )Nr3   Trw   rE   s    r   atomic_add_double_global_2_wraprz      s    c1dkkoo7GNr   c                 Z    [        U S[        R                  R                  [        S5        g rv   )r<   r   rC   rD   r   rE   s    r   atomic_add_double_global_3r|      s    c1dkkoo7L#%r   c                 h    [        XS[        S[        R                  R                  [
        SS5	        g rX   r.   r   r   rC   rD   r   rp   s     r   atomic_add_doubler      %    cWb"kkoo/?eMr   c                 h    [        XS[        S[        R                  R                  [
        SS5	        g r_   r~   rp   s     r   atomic_add_double_wrapr      s%    cWb"kkoo/?dLr   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rb   r9   r   r   rC   rD   r   rE   s    r   atomic_add_double_2r      re   r   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rg   r   rE   s    r   atomic_add_double_2_wrapr      ri   r   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rb   )r9   r   r   rC   rD   r   rE   s    r   atomic_add_double_3r      rl   r   c                 h    [        X S[        S[        R                  R                  [
        SS5	        g r@   )r.   r   r   rC   subr   rE   s    r   
atomic_subr      rG   r   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rK   )r9   r   r   rC   r   r   rE   s    r   atomic_sub2r      rQ   r   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rK   )r9   r   r   rC   r   r   rE   s    r   atomic_sub3r      rV   r   c                 h    [        X S[        S[        R                  R                  [
        SS5	        g rX   )r.   r   r   rC   r   r   rE   s    r   atomic_sub_floatr      r]   r   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rb   )r9   r   r   rC   r   r   rE   s    r   atomic_sub_float_2r      re   r   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rb   )r9   r   r   rC   r   r   rE   s    r   atomic_sub_float_3r      rl   r   c                 h    [        XS[        S[        R                  R                  [
        SS5	        g rX   )r.   r   r   rC   r   r   rp   s     r   atomic_sub_doubler      r   r   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rb   )r9   r   r   rC   r   r   rE   s    r   atomic_sub_double_2r      re   r   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rb   r9   r   r   rC   r   r   rE   s    r   atomic_sub_double_3r      rl   r   c                 R    [        XSS[        R                  R                  S5        g rn   )r>   r   rC   r   rp   s     r   atomic_sub_double_globalr      rr   r   c                 Z    [        U S[        R                  R                  [        S5        g )NrY   F)r<   r   rC   r   r   rE   s    r   atomic_sub_double_global_2r      s    c39I#%r   c           	      f    [        U S[        S[        R                  R                  [
        S5        g rb   r   rE   s    r   atomic_sub_double_global_3r      rl   r   c                 h    [        X U[        S[        R                  R                  [
        SS5	        g )NrA   r3   F)r.   r   r   rC   and_r   r"   r$   s     r   
atomic_andr      s'    cVR"kk..0@!ULr   c           	      d    [        X[        S[        R                  R                  [
        S5        g NrL   F)r9   r   r   rC   r   r   r   s     r   atomic_and2r      #    c"kk..0@%Ir   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   atomic_and3r      s#    c"kk..0EuNr   c                 R    [        XSU[        R                  R                  S5        g NrA   F)r>   r   rC   r   r#   r"   r$   s      r   atomic_and_globalr         cC1A1A5Ir   c                 X    [        X[        R                  R                  [        S5        g NF)r<   r   rC   r   r   r   s     r   atomic_and_global_2r     s    c(8(8.7r   c                 h    [        X U[        S[        R                  R                  [
        SS5	        g NrA   r   F)r.   r   r   rC   or_r   r   s     r   	atomic_orr     %    cVR"kkoo/?EKr   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   
atomic_or2r     !    c"kkoo/?Hr   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   
atomic_or3r     !    c"kkoo/DeMr   c                 R    [        XSU[        R                  R                  S5        g r   )r>   r   rC   r   r   s      r   atomic_or_globalr     rr   r   c                 X    [        X[        R                  R                  [        S5        g r   )r<   r   rC   r   r   r   s     r   atomic_or_global_2r         c.7r   c                 h    [        X U[        S[        R                  R                  [
        SS5	        g r   )r.   r   r   rC   xorr   r   s     r   
atomic_xorr   $  r   r   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   atomic_xor2r   )  r   r   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   atomic_xor3r   .  r   r   c                 R    [        XSU[        R                  R                  S5        g r   )r>   r   rC   r   r   s      r   atomic_xor_globalr   3  rr   r   c                 X    [        X[        R                  R                  [        S5        g r   )r<   r   rC   r   r   r   s     r   atomic_xor_global_2r   7  r   r   c           	      d    [        XU[        S[        R                  R                  [
        5        g NrA   )r0   r   r   rC   incr   r"   r#   r$   s      r   atomic_inc32r   <  !    sfb#{{0@Br   c           	      d    [        XU[        S[        R                  R                  [
        5        g r   )r0   r   r   rC   r   r   r   s      r   atomic_inc64r   A  !    sfb#{{0BDr   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   atomic_inc2_32r   F  r   r   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   atomic_inc2_64r   K  r   r   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   atomic_inc3r   P  r   r   c                 R    [        XSU[        R                  R                  S5        g r   )r>   r   rC   r   r   s      r   atomic_inc_globalr   U  rr   r   c                 X    [        X[        R                  R                  [        S5        g r   )r<   r   rC   r   r   r   s     r   atomic_inc_global_2r   Y  r   r   c           	      d    [        XU[        S[        R                  R                  [
        5        g r   )r0   r   r   rC   decr   r   s      r   atomic_dec32r   ^  r   r   c           	      d    [        XU[        S[        R                  R                  [
        5        g r   )r0   r   r   rC   r   r   r   s      r   atomic_dec64r   c  r   r   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   atomic_dec2_32r   h  r   r   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   atomic_dec2_64r   m  r   r   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   atomic_dec3r   r  r   r   c                 R    [        XSU[        R                  R                  S5        g r   )r>   r   rC   r   r   s      r   atomic_dec_globalr   w  rr   r   c                 X    [        X[        R                  R                  [        S5        g r   )r<   r   rC   r   r   r   s     r   atomic_dec_global_2r   {  r   r   c           	      d    [        XU[        S[        R                  R                  [
        5        g r   )r0   r   r   rC   exchr   r   s      r   atomic_exchr     s#    sfb#{{//1ACr   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   atomic_exch2r     r   r   c           	      d    [        X[        S[        R                  R                  [
        S5        g r   )r9   r   r   rC   r   r   r   s     r   atomic_exch3r     r   r   c                 R    [        XSU[        R                  R                  S5        g r   )r>   r   rC   r   r   s      r   atomic_exch_globalr     r   r   c                     [        S5      R                  U S9n0 n[        U[        [        [
        S.U5        US   US   US   US   4$ )Na  
    def atomic(res, ary):
        tx = cuda.threadIdx.x
        bx = cuda.blockIdx.x
        {func}(res, 0, ary[tx, bx])

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

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

    def atomic_double_shared(res, ary):
        tid = cuda.threadIdx.x
        smary = cuda.shared.array(32, float64)
        smary[tid] = ary[tid]
        smres = cuda.shared.array(1, float64)
        if tid == 0:
            smres[0] = res[0]
        cuda.syncthreads()
        {func}(smres, 0, smary[tid])
        cuda.syncthreads()
        if tid == 0:
            res[0] = smres[0]
    )func)r   r   r   rC   atomic_double_normalizedindexatomic_double_oneindexatomic_double_shared)r   formatexecr   r   r   )r   fnslds      r   gen_atomic_extreme_funcsr    sh    
  	6 
T	7 8 
Bt6BBGxL"<='("-C*DF Fr   zcuda.atomic.maxzcuda.atomic.minzcuda.atomic.nanmaxzcuda.atomic.nanminc                     [         R                  " S5      nX@R                  :  a(  [         R                  R	                  XS  X2U   5      X'   g g Nr3   )r   gridsizerC   compare_and_swapresoldr"   fill_valgids        r   atomic_compare_and_swapr    s=    
))A,C
XX~;;//D	8XN r   c                     [         R                  " S5      nX@R                  :  a&  [         R                  R	                  XX2U   5      X'   g g r  )r   r  r	  rC   casr  s        r   atomic_cas_1dimr    s7    
))A,C
XX~;;??3X3x@ r   c                     [         R                  " S5      nUS   U R                  S   :  a=  US   U R                  S   :  a&  [         R                  R	                  XX2U   5      X'   g g g )N   r   r3   )r   r  r;   rC   r  r  s        r   atomic_cas_2dimr    sX    
))A,C
1v		!Q#))A,!6;;??3X3x@ "7r   c                     ^  \ rS rSrU 4S jrS rS rS rS rS r	S r
SS	 jrS
 rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS r S r!S  r"S! r#S" r$S# r%S$ r&S% r'S& r(S' r)S( r*S) r+S* r,S+ r-S, r.S- r/S. r0S/ r1S0 r2S1 r3S2 r4S3 r5S4 r6S5 r7S6 r8S7 r9S8 r:S9 r;S: r<S; r=S< r>S= r?S> r@S? rAS@ rBSA rCSB rDSC rESD rFSE rGSF rHSG rISH rJSI rKSJ rLSK rMSL rNSM rOSN rPSO rQSP rRSQ rSSR rTSS rUST rVSU rWSV rXSW rYSX rZSY r[SZ r\S[ r]S\ r^S] r_S^ r`S_ raS` rbSa rcSb rdSSc jreSd rfSe rgSf rhSg riSh rjSi rkSj rlSk rmSl rnSm roSn rpSo rqSp rrSq rsSr rtSs ruSt rvSu rwSv rxSw rySx rzSy r{Sz r|S{ r}S| r~S} rS~ rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rSrU =r$ )TestCudaAtomicsi  c                 `   > [         TU ]  5         [        R                  R	                  S5        g )Nr   )supersetUpnprandomseed)self	__class__s    r   r  TestCudaAtomics.setUp  s    
		qr   c                    [         R                  R                  SSSS9R                  [         R                  5      nUR                  5       nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        [        R                  " S5      " [        5      nUS   " U5        [         R                  " S[         R                  S9n[        UR                  5       H  nXcU   ==   S-  ss'   M     U R                  [         R                  " X:H  5      5        U R                  [         R                  " X&:H  5      5        g Nr   rA   r	  zvoid(uint32[:])r3   rA   dtyper3   )r  r  randintastyper   copyr   jitrF   rI   zerosranger	  
assertTrueall)r  r"   ary_wraporigcuda_atomic_addcuda_atomic_add_wrapgoldis           r   test_atomic_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                 0   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        [        R                  " S5      " [        5      nUS   " U5        U R                  [         R                  " XS-   :H  5      5        U R                  [         R                  " X#S-   :H  5      5        g 	Nr   rA   r$  rM   rN   zvoid(uint32[:,:])r3   rL   r3   )r  r  r(  r)  r   reshaper*  r   r+  rP   rS   r.  r/  )r  r"   r0  r1  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                 f   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        U R                  [         R                  " XS-   :H  5      5        g r9  )r  r  r(  r)  r   r;  r*  r   r+  rU   r.  r/  r  r"   r1  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                    [         R                  R                  SSSS9R                  [         R                  5      nUR                  5       nUR                  5       R                  [         R                  5      n[        R                  " S5      " [        5      nUS   " U5        [        R                  " S5      " [        5      nUS   " U5        [         R                  " S[         R                  S9n[        UR                  5       H  nXcU   ==   S-  ss'   M     U R                  [         R                   " X:H  5      5        U R                  [         R                   " X&:H  5      5        g Nr   rA   r$  zvoid(float32[:])r%  r&  rY   )r  r  r(  r)  r   r*  intpr   r+  r\   r`   r,  r   r-  r	  r.  r/  )r  r"   r0  r1  cuda_atomic_add_floatadd_float_wrapr4  r5  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                 0   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        [        R                  " S5      " [        5      nUS   " U5        U R                  [         R                  " XS-   :H  5      5        U R                  [         R                  " X#S-   :H  5      5        g 	Nr   rA   r$  rM   rN   zvoid(float32[:,:])r:  r3   )r  r  r(  r)  r   r;  r*  r   r+  rd   rh   r.  r/  )r  r"   r0  r1  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                 f   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        U R                  [         R                  " XS-   :H  5      5        g rM  )r  r  r(  r)  r   r;  r*  r   r+  rk   r.  r/  rA  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                    [         R                  (       a  g [        [        UR	                  5       R                  5       5      5      n[        SS5      (       aI  [        R                  R                  5       S:  a  SnOSnU(       a  U S3nU R                  U S3U5        g U(       a  U R                  SU5        g U R                  S	U5        g )
N   r   )   r3   redatomz.sharedz.add.f64zatom.shared.cas.b64zatom.cas.b64)r   ENABLE_CUDASIMnextiterinspect_asmvaluesr   r   runtimeget_versionassertIn)r  kernelr   asminsts        r   assertCorrectFloat64Atomics+TestCudaAtomics.assertCorrectFloat64Atomics*  s       4**,33567A||'')G3 w'MMTF(+S13S9nc2r   c                    [         R                  R                  SSS[         R                  S9n[         R                  " S[         R
                  5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X5        [        R                  " S5      " [        5      nUS   " X5        [         R                  " S[         R                  S9n[        UR                  5       H  nXaU   ==   S-  ss'   M     [         R                  R                  X&5        [         R                  R                  X65        U R!                  U5        U R!                  U5        g Nr   rA   r	  r'  void(int64[:], float64[:])r%  r&  rY   )r  r  r(  int64r,  r   r*  r   r+  r   r   r   r-  r	  testingassert_equalrd  )r  r#   r"   r0  cuda_fnwrap_fnr4  r5  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                 T   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        [        R                  " S5      " [        5      nUS   " U5        [         R                  R                  XS-   5        [         R                  R                  X#S-   5        U R                  U5        U R                  U5        g 	Nr   rA   r$  rM   rN   void(float64[:,:])r:  r3   )r  r  r(  r)  r   r;  r*  r   r+  r   r   rk  rl  rd  )r  r"   r0  r1  rm  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                 x   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        [         R                  R                  XS-   5        U R                  U5        g rr  )r  r  r(  r)  r   r;  r*  r   r+  r   rk  rl  rd  r  r"   r1  	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                    [         R                  R                  SSS[         R                  S9n[         R                  " S[         R
                  5      nUR                  5       nSn[        R                  " U5      " [        5      n[        R                  " U5      " [        5      nUS   " X5        US   " X5        [         R                  " S[         R                  S9n[        UR                  5       H  nXqU   ==   S-  ss'   M     [         R                  R                  X'5        [         R                  R                  X75        U R!                  USS	9  U R!                  USS	9  g )
Nr   rA   rh  ri  r%  r&  rY   Fr   )r  r  r(  rj  r,  r   r*  r   r+  rq   rt   r   r-  r	  rk  rl  rd  )	r  r#   r"   r0  sigry  wrap_cuda_funcr4  r5  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                 T   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       nUR                  5       nSn[        R                  " U5      " [        5      n[        R                  " U5      " [        5      nUS   " U5        US   " U5        [         R                  R                  XS-   5        [         R                  R                  X#S-   5        U R                  US	S
9  U R                  US	S
9  g Nr   rA   r$  rM   rN   rs  r:  r3   Fr}  )r  r  r(  r)  r   r;  r*  r   r+  rx   rz   rk  rl  rd  )r  r"   r0  r1  r~  ry  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                 v   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        [         R                  R                  XS-   5        U R                  US	S
9  g r  )r  r  r(  r)  r   r;  r*  r   r+  r|   rk  rl  rd  rx  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                    [         R                  R                  SSSS9R                  [         R                  5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        [         R                  " S[         R                  S9n[        UR                  5       H  nXBU   ==   S-  ss'   M     U R                  [         R                  " X:H  5      5        g r#  )r  r  r(  r)  r   r*  r   r+  r   r,  r-  r	  r.  r/  )r  r"   r1  cuda_atomic_subr4  r5  s         r   test_atomic_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                 f   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        U R                  [         R                  " XS-
  :H  5      5        g r9  )r  r  r(  r)  r   r;  r*  r   r+  r   r.  r/  r  r"   r1  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                 f   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        U R                  [         R                  " XS-
  :H  5      5        g r9  )r  r  r(  r)  r   r;  r*  r   r+  r   r.  r/  r  r"   r1  cuda_atomic_sub3s       r   test_atomic_sub3 TestCudaAtomics.test_atomic_sub3  r  r   c                    [         R                  R                  SSSS9R                  [         R                  5      nUR                  5       R                  [         R                  5      n[        R                  " S5      " [        5      nUS   " U5        [         R                  " S[         R                  S9n[        UR                  5       H  nXBU   ==   S-  ss'   M     U R                  [         R                  " X:H  5      5        g rF  )r  r  r(  r)  r   r*  rG  r   r+  r   r,  r-  r	  r.  r/  )r  r"   r1  cuda_atomic_sub_floatr4  r5  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                 f   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        U R                  [         R                  " XS-
  :H  5      5        g rM  )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                 f   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        U R                  [         R                  " XS-
  :H  5      5        g rM  )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                    [         R                  R                  SSS[         R                  S9n[         R                  " S[         R
                  5      n[        R                  " S5      " [        5      nUS   " X5        [         R                  " S[         R
                  S9n[        UR                  5       H  nXAU   ==   S-  ss'   M     [         R                  R                  X$5        g rg  )r  r  r(  rj  r,  r   r   r+  r   r-  r	  rk  rl  )r  r#   r"   ry  r4  r5  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                 V   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        [         R                  R                  XS-
  5        g rr  )r  r  r(  r)  r   r;  r*  r   r+  r   rk  rl  rx  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                 V   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        [         R                  R                  XS-
  5        g rr  )r  r  r(  r)  r   r;  r*  r   r+  r   rk  rl  rx  s       r   test_atomic_sub_double_3(TestCudaAtomics.test_atomic_sub_double_3  r  r   c                    [         R                  R                  SSS[         R                  S9n[         R                  " S[         R
                  5      nSn[        R                  " U5      " [        5      nUS   " X5        [         R                  " S[         R
                  S9n[        UR                  5       H  nXQU   ==   S-  ss'   M     [         R                  R                  X%5        g rg  )r  r  r(  rj  r,  r   r   r+  r   r-  r	  rk  rl  )r  r#   r"   r~  ry  r4  r5  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                 V   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        [         R                  R                  XS-
  5        g rr  )r  r  r(  r)  r   r;  r*  r   r+  r   rk  rl  rx  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                 V   [         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " U5        [         R                  R                  XS-
  5        g rr  )r  r  r(  r)  r   r;  r*  r   r+  r   rk  rl  rx  s       r   test_atomic_sub_double_global_3/TestCudaAtomics.test_atomic_sub_double_global_3  r  r   c                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        UR                  5       n[        UR                  5       H  nXSU   ==   U-  ss'   M     U R                  [         R                  " X%:H  5      5        g )N  r   rA   r$  void(uint32[:], uint32)r%  )r  r  r(  r)  r   r*  r   r+  r   r-  r	  r.  r/  r  
rand_constr"   r1  ry  r4  r5  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                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        U R                  [         R                  " X#U-  :H  5      5        g 	Nr  r   rA   r$  rM   rN   void(uint32[:,:], uint32)r:  )r  r  r(  r)  r   r;  r*  r   r+  r   r.  r/  r  r  r"   r1  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                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        U R                  [         R                  " X#U-  :H  5      5        g r  )r  r  r(  r)  r   r;  r*  r   r+  r   r.  r/  r  r  r"   r1  cuda_atomic_and3s        r   test_atomic_and3 TestCudaAtomics.test_atomic_and3  r  r   c                 
   [         R                  R                  S5      n[         R                  R                  SSS[         R                  S9n[         R                  R                  SSS[         R                  S9nSn[        R
                  " U5      " [        5      nUS   " X#U5        UR                  5       n[        UR                  5       H  nXbU   ==   U-  ss'   M     [         R                  R                  X65        g Nr  r   rA   rh  zvoid(int32[:], int32[:], int32)r%  )r  r  r(  int32r   r+  r   r*  r-  r	  rk  rl  r  r  r#   r"   r~  ry  r4  r5  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                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        [         R                  R                  X#U-  5        g r  )r  r  r(  r)  r   r;  r*  r   r+  r   rk  rl  r  r  r"   r1  ry  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                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        [         R                  " S[         R                  S9n[        UR                  5       H  nXSU   ==   U-  ss'   M     U R                  [         R                  " X%:H  5      5        g Nr  r   rA   r$  r  r%  r&  )r  r  r(  r)  r   r*  r   r+  r   r,  r-  r	  r.  r/  r  s          r   test_atomic_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                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        U R                  [         R                  " X#U-  :H  5      5        g 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                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        U R                  [         R                  " X#U-  :H  5      5        g 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                 
   [         R                  R                  S5      n[         R                  R                  SSS[         R                  S9n[         R                  R                  SSS[         R                  S9nSn[        R
                  " U5      " [        5      nUS   " X#U5        UR                  5       n[        UR                  5       H  nXbU   ==   U-  ss'   M     [         R                  R                  X65        g r  )r  r  r(  r  r   r+  r   r*  r-  r	  rk  rl  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                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        [         R                  R                  X#U-  5        g r  )r  r  r(  r)  r   r;  r*  r   r+  r   rk  rl  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                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        [         R                  " S[         R                  S9n[        UR                  5       H  nXSU   ==   U-  ss'   M     U R                  [         R                  " X%:H  5      5        g 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                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        U R                  [         R                  " X#U-  :H  5      5        g r  )r  r  r(  r)  r   r;  r*  r   r+  r   r.  r/  )r  r  r"   r1  cuda_atomic_xor2s        r   test_atomic_xor2 TestCudaAtomics.test_atomic_xor2{  r  r   c                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        U R                  [         R                  " X#U-  :H  5      5        g r  )r  r  r(  r)  r   r;  r*  r   r+  r   r.  r/  )r  r  r"   r1  cuda_atomic_xor3s        r   test_atomic_xor3 TestCudaAtomics.test_atomic_xor3  r  r   c                 
   [         R                  R                  S5      n[         R                  R                  SSS[         R                  S9n[         R                  R                  SSS[         R                  S9nUR	                  5       nSn[
        R                  " U5      " [        5      nUS   " X#U5        [        UR                  5       H  nXBU   ==   U-  ss'   M     [         R                  R                  X45        g r  )r  r  r(  r  r*  r   r+  r   r-  r	  rk  rl  )r  r  r#   r"   r4  r~  ry  r5  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                    [         R                  R                  S5      n[         R                  R                  SSSS9R                  [         R                  5      R                  SS5      nUR                  5       n[        R                  " S5      " [        5      nUS   " X!5        [         R                  R                  X#U-  5        g r  )r  r  r(  r)  r   r;  r*  r   r+  r   rk  rl  r  s        r   test_atomic_xor_global_2(TestCudaAtomics.test_atomic_xor_global_2  r  r   c                     [         R                  R                  SUS9n[         R                  R                  SSSS9R                  U5      n[         R                  " SUS9nX#U4$ )NrA   r&  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                     [         R                  R                  SUS9n[         R                  R                  SSSS9R                  U5      R	                  SS5      nX#4$ )NrA   r&  r   r$  rM   rN   )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           	          UR                  5       n[        R                  " U5      " U5      n	XU4   " XU5        [        R                  R                  U[        R                  " X:  SUS-   5      5        g r2   r*  r   r+  r  rk  rl  where
r  r"   r#   r  r~  nblocksblksizer   r1  ry  s
             r   check_inc_indexTestCudaAtomics.check_inc_index  X    xxzHHSM$'	7"#Cf5


RXXdna%JKr   c           	          UR                  5       n[        R                  " U5      " U5      n	XU4   " X!U5        [        R                  R                  U[        R                  " X:  SUS-   5      5        g r2   r  r  s
             r   check_inc_index2 TestCudaAtomics.check_inc_index2  r
  r   c           	          UR                  5       n[        R                  " U5      " U5      nXU4   " X5        [        R                  R                  U[        R                  " Xr:  SUS-   5      5        g r2   r  	r  r"   r  r~  r  r  r   r1  ry  s	            r   	check_incTestCudaAtomics.check_inc  sV    xxzHHSM$'	7"#C0


RXXdna%JKr   c           	      z    U R                  [        R                  S9u  pnSnU R                  X#XSS[        5        g Nr&  "void(uint32[:], uint32[:], uint32)r3   rA   )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           	      z    U R                  [        R                  S9u  pnSnU R                  X#XSS[        5        g Nr&  z"void(uint64[:], uint64[:], uint64)r3   rA   )r  r  r   r  r   r  s        r   test_atomic_inc_64"TestCudaAtomics.test_atomic_inc_64  r  r   c                 |    U R                  [        R                  5      u  pSnU R                  X!USS[        5        g Nr  r3   rL   )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                 |    U R                  [        R                  5      u  pSnU R                  X!USS[        5        g Nvoid(uint64[:,:], uint64)r3   rL   )r   r  r   r  r   r  s       r   test_atomic_inc2_64#TestCudaAtomics.test_atomic_inc2_64  r"  r   c                 |    U R                  [        R                  5      u  pSnU R                  X!USS[        5        g r  )r   r  r   r  r   r  s       r   test_atomic_inc3 TestCudaAtomics.test_atomic_inc3  1    11"))<
)sQ{Cr   c           	      z    U R                  [        R                  S9u  pnSnU R                  X#XSS[        5        g 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           	      z    U R                  [        R                  S9u  pnSnU R                  X#XSS[        5        g r  )r  r  r   r  r   r  s        r   test_atomic_inc_global_64)TestCudaAtomics.test_atomic_inc_global_64  r/  r   c                 |    U R                  [        R                  5      u  pSnU R                  X!USS[        5        g 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                 |    U R                  [        R                  5      u  pSnU R                  X!USS[        5        g r$  )r   r  r   r  r   r  s       r   test_atomic_inc_global_2_64+TestCudaAtomics.test_atomic_inc_global_2_64  r6  r   c                    UR                  5       n[        R                  " U5      " U5      n	XU4   " XU5        [        R                  R                  U[        R                  " US:H  U[        R                  " X:  UUS-
  5      5      5        g 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                    UR                  5       n[        R                  " U5      " U5      n	XU4   " X!U5        [        R                  R                  U[        R                  " US:H  U[        R                  " X:  UUS-
  5      5      5        g r2   r  r  s
             r   check_dec_index2 TestCudaAtomics.check_dec_index2  r=  r   c                    UR                  5       n[        R                  " U5      " U5      nXU4   " X5        [        R                  R                  U[        R                  " US:H  U[        R                  " Xr:  UUS-
  5      5      5        g 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           	      z    U R                  [        R                  S9u  pnSnU R                  X#XSS[        5        g r  )r  r  r   r;  r   r  s        r   test_atomic_dec_32"TestCudaAtomics.test_atomic_dec_32  r  r   c           	      z    U R                  [        R                  S9u  pnSnU R                  X#XSS[        5        g r  )r  r  r   r;  r   r  s        r   test_atomic_dec_64"TestCudaAtomics.test_atomic_dec_64  r  r   c                 |    U R                  [        R                  5      u  pSnU R                  X!USS[        5        g r  )r   r  r   rB  r   r  s       r   test_atomic_dec2_32#TestCudaAtomics.test_atomic_dec2_32  r"  r   c                 |    U R                  [        R                  5      u  pSnU R                  X!USS[        5        g r$  )r   r  r   rB  r   r  s       r   test_atomic_dec2_64#TestCudaAtomics.test_atomic_dec2_64  r"  r   c                 |    U R                  [        R                  5      u  pSnU R                  X!USS[        5        g r  )r   r  r   rB  r   r  s       r   test_atomic_dec3_new$TestCudaAtomics.test_atomic_dec3_new  r+  r   c           	      z    U R                  [        R                  S9u  pnSnU R                  X#XSS[        5        g r  )r  r  r   r?  r   r  s        r   test_atomic_dec_global_32)TestCudaAtomics.test_atomic_dec_global_32!  r/  r   c           	      z    U R                  [        R                  S9u  pnSnU R                  X#XSS[        5        g r  )r  r  r   r?  r   r  s        r   test_atomic_dec_global_64)TestCudaAtomics.test_atomic_dec_global_64'  r/  r   c                 |    U R                  [        R                  5      u  pSnU R                  X!USS[        5        g r  )r   r  r   rB  r   r  s       r   test_atomic_dec_global2_32*TestCudaAtomics.test_atomic_dec_global2_32-  r6  r   c                 |    U R                  [        R                  5      u  pSnU R                  X!USS[        5        g r$  )r   r  r   rB  r   r  s       r   test_atomic_dec_global2_64*TestCudaAtomics.test_atomic_dec_global2_642  r6  r   c                    [         R                  R                  SS[         R                  S9n[         R                  R                  SSSS9R	                  [         R                  5      n[         R
                  " S[         R                  S9n[        R                  " S5      " [        5      nUS   " X#U5        [         R                  R                  X!5        g )	N2   d   r&  r   rA   r$  r  r%  )r  r  r(  r   r)  r  r   r+  r   rk  rl  )r  r  r"   r#   ry  s        r   test_atomic_exch TestCudaAtomics.test_atomic_exch7  s    YY&&r3bii&@
ii2B/66ryyAii")),HHAB;O	%:.


0r   c                    [         R                  R                  SS[         R                  S9n[         R                  R                  SSSS9R	                  [         R                  5      R                  SS5      n[        R                  " S	5      " [        5      nUS
   " X!5        [         R                  R                  X!5        g )Nr`  ra  r&  r   rA   r$  rM   rN   r  r:  )r  r  r(  r   r)  r;  r   r+  r   rk  rl  r  r  r"   ry  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                    [         R                  R                  SS[         R                  S9n[         R                  R                  SSSS9R	                  [         R                  5      R                  SS5      n[        R                  " S	5      " [        5      nUS
   " X!5        [         R                  R                  X!5        g )Nr`  ra  r&  r   rA   r$  rM   rN   r%  r:  )r  r  r(  r   r)  r;  r   r+  r   rk  rl  re  s       r   test_atomic_exch3!TestCudaAtomics.test_atomic_exch3I  rh  r   c                    [         R                  R                  SS[         R                  S9n[         R                  " S[         R                  S9n[         R                  R                  SSS[         R                  S9nSn[
        R                  " U5      " [        5      nUS   " X#U5        [         R                  R                  X15        g )	Nr`  ra  r&  rA   r   rh  r  r%  )
r  r  r(  r   r  r   r+  r   rk  rl  )r  r  r#   r"   r~  ry  s         r   test_atomic_exch_global'TestCudaAtomics.test_atomic_exch_globalQ  s    YY&&r3bii&@
ii")),ii2Bbii@2HHSM"45	%:.


0r   c                 P   [         R                  R                  X#SS9R                  U5      n[         R                  " SUR
                  S9n[        R                  " [        5      nUS   " XT5        [         R                  " U5      n[         R                  R                  XW5        g )NrA   rA   r$  r3   r&  )r  r  r(  r)  r,  r'  r   r+  
atomic_maxmaxrk  rl  r  r'  lohivalsr  ry  r4  s           r   check_atomic_max TestCudaAtomics.check_atomic_max[  st    yy  h 7>>uEhhq

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


*r   c                 B    U R                  [        R                  SSS9  g N   r'  rt  ru  )rw  r  r  r  s    r   test_atomic_max_int32%TestCudaAtomics.test_atomic_max_int32c      BHHEBr   c                 B    U R                  [        R                  SSS9  g Nr   r|  r}  )rw  r  r   r~  s    r   test_atomic_max_uint32&TestCudaAtomics.test_atomic_max_uint32f      BII!>r   c                 B    U R                  [        R                  SSS9  g rz  )rw  r  rj  r~  s    r   test_atomic_max_int64%TestCudaAtomics.test_atomic_max_int64i  r  r   c                 B    U R                  [        R                  SSS9  g r  )rw  r  r   r~  s    r   test_atomic_max_uint64&TestCudaAtomics.test_atomic_max_uint64l  r  r   c                 B    U R                  [        R                  SSS9  g rz  )rw  r  r   r~  s    r   test_atomic_max_float32'TestCudaAtomics.test_atomic_max_float32o      BJJ6eDr   c                 B    U R                  [        R                  SSS9  g rz  )rw  r  r   r~  s    r   test_atomic_max_double&TestCudaAtomics.test_atomic_max_doubler  r  r   c                    [         R                  R                  SSSS9R                  [         R                  5      n[         R
                  " S[         R                  5      n[        R                  " S5      " [        5      nUS   " X!5        [         R                  " U5      n[         R                  R                  X$5        g Nr   r|  rp  r$  r3   void(float64[:], float64[:,:]))r  r  r(  r)  r   r,  r   r+  !atomic_max_double_normalizedindexrr  rk  rl  r  rv  r  ry  r4  s        r   &test_atomic_max_double_normalizedindex6TestCudaAtomics.test_atomic_max_double_normalizedindexu  s    yy  E 9@@Lhhq"**%HH=>-/	&#$vvd|


*r   c                    [         R                  R                  SSSS9R                  [         R                  5      n[         R
                  " S[         R                  5      n[        R                  " S5      " [        5      nUS   " X!5        [         R                  " U5      n[         R                  R                  X$5        g Nr      rA   r$  r3   void(float64[:], float64[:])r%  )r  r  r(  r)  r   r,  r   r+  atomic_max_double_oneindexrr  rk  rl  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                 R   [         R                  R                  X#SS9R                  U5      n[         R                  " S/UR
                  S9n[        R                  " [        5      nUS   " XT5        [         R                  " U5      n[         R                  R                  XW5        g )Nrp  r$  r|  r&  )r  r  r(  r)  r    r'  r   r+  
atomic_minminrk  rl  rs  s           r   check_atomic_min TestCudaAtomics.check_atomic_min  sv    yy  h 7>>uEhhwdjj1HHZ(	&#$vvd|


*r   c                 B    U R                  [        R                  SSS9  g rz  )r  r  r  r~  s    r   test_atomic_min_int32%TestCudaAtomics.test_atomic_min_int32  r  r   c                 B    U R                  [        R                  SSS9  g r  )r  r  r   r~  s    r   test_atomic_min_uint32&TestCudaAtomics.test_atomic_min_uint32  r  r   c                 B    U R                  [        R                  SSS9  g rz  )r  r  rj  r~  s    r   test_atomic_min_int64%TestCudaAtomics.test_atomic_min_int64  r  r   c                 B    U R                  [        R                  SSS9  g r  )r  r  r   r~  s    r   test_atomic_min_uint64&TestCudaAtomics.test_atomic_min_uint64  r  r   c                 B    U R                  [        R                  SSS9  g rz  )r  r  r   r~  s    r   test_atomic_min_float%TestCudaAtomics.test_atomic_min_float  r  r   c                 B    U R                  [        R                  SSS9  g rz  )r  r  r   r~  s    r   test_atomic_min_double&TestCudaAtomics.test_atomic_min_double  r  r   c                    [         R                  R                  SSSS9R                  [         R                  5      n[         R
                  " S[         R                  5      S-  n[        R                  " S5      " [        5      nUS   " X!5        [         R                  " U5      n[         R                  R                  X$5        g r  )r  r  r(  r)  r   onesr   r+  !atomic_min_double_normalizedindexr  rk  rl  r  s        r   &test_atomic_min_double_normalizedindex6TestCudaAtomics.test_atomic_min_double_normalizedindex  s    yy  E 9@@Lgga$u,HH=>-/	&#$vvd|


*r   c                    [         R                  R                  SSSS9R                  [         R                  5      n[         R
                  " S[         R                  5      S-  n[        R                  " S5      " [        5      nUS   " X!5        [         R                  " U5      n[         R                  R                  X$5        g r  )r  r  r(  r)  r   r  r   r+  atomic_min_double_oneindexr  rk  rl  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                    [         R                  " S5      " U5      n[        R                  R	                  SSSS9R                  [        R                  5      n[        R                  " S[        R                  5      [        R                  -   nUS   " XC5        [        R                  R                  U[        R                  /5        g )Nr  r   r  r3   r3   r$  r3   )r   r+  r  r  r(  r)  r   r,  nanrk  rl  )r  r   ry  rv  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                    [         R                  " S5      " U5      n[        R                  R	                  SSSS9R                  [        R                  5      nUR                  5       n[        R                  " S[        R                  5      [        R                  -   nUS   " X55        [        R                  R                  X45        g )Nr  r   r  r3   r$  r  )r   r+  r  r  r(  r)  r   r*  r,  r  rk  rl  )r  r   ry  r  r4  rv  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                 .    U R                  [        5        g r   )r  r  r~  s    r   test_atomic_min_nan_location,TestCudaAtomics.test_atomic_min_nan_location      --j9r   c                 .    U R                  [        5        g r   )r  rq  r~  s    r   test_atomic_max_nan_location,TestCudaAtomics.test_atomic_max_nan_location  r  r   c                 .    U R                  [        5        g r   )r  r  r~  s    r   test_atomic_min_nan_val'TestCudaAtomics.test_atomic_min_nan_val      ((4r   c                 .    U R                  [        5        g r   )r  rq  r~  s    r   test_atomic_max_nan_val'TestCudaAtomics.test_atomic_max_nan_val  r  r   c                    [         R                  R                  SSSS9R                  [         R                  5      n[         R
                  " S[         R                  5      nSn[        R                  " U5      " [        5      nUS   " X!5        [         R                  " U5      n[         R                  R                  X%5        g Nr   rA   r$  r3   r  r%  )r  r  r(  r)  r   r,  r   r+  atomic_max_double_sharedrr  rk  rl  r  rv  r  r~  ry  r4  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                    [         R                  R                  SSSS9R                  [         R                  5      n[         R
                  " S[         R                  5      S-  nSn[        R                  " U5      " [        5      nUS   " X!5        [         R                  " U5      n[         R                  R                  X%5        g r  )r  r  r(  r)  r   r  r   r+  atomic_min_double_sharedr  rk  rl  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                    U/US-  -  U/US-  -  -   n[         R                  R                  U5        [         R                  " XtS9nUS:X  a  SUl        [         R
                  " U5      n[         R                  R                  SSUR                  S9R                  UR                  5      n	Xr:H  n
Xs:H  n[         R
                  " U5      nX   X'   X<U'   UR                  5       n[        R                  " U5      nUS:X  a  US   " XxX5        OUS   " XxX5        [         R                  R                  X5        [         R                  R                  X5        g )	Nr  r&  )
   r3   r  r$  r  r  )r  r  )r  r  shuffleasarrayr;   
zeros_liker(  r)  r'  r*  r   r+  rk  assert_array_equal)r  nfillunfillr'  cas_funcndimr  outr"   	fill_maskunfill_mask
expect_res
expect_outry  s                  r   	check_casTestCudaAtomics.check_cas  s   fQ6(a1f"55
		#jj*19 CImmC ii2CII6==ciiHK	m]]3'
 #
"(;XXZ
HHX&	19fc2()#C>


%%j6


%%j6r   c                 N    U R                  SSS[        R                  [        S9  g Nra  r  r  r  r  r'  r  )r  r  r  r  r~  s    r   test_atomic_compare_and_swap,TestCudaAtomics.test_atomic_compare_and_swap  "    3r 7 	 	9r   c                 N    U R                  SSS[        R                  [        S9  g Nra  r  r  )r  r  rj  r  r~  s    r   test_atomic_compare_and_swap2-TestCudaAtomics.test_atomic_compare_and_swap2  r  r   c                     [         R                  R                  SS[         R                  S9n[         R                  R                  SS[         R                  S9nU R	                  SX[         R                  [
        S9  g Nr`  r  r&  r3      ra  r  )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                     [         R                  R                  SS[         R                  S9n[         R                  R                  SS[         R                  S9nU R	                  SX[         R                  [
        S9  g r  )r  r  r(  r   r  r  r  s      r   test_atomic_compare_and_swap4-TestCudaAtomics.test_atomic_compare_and_swap4  r
  r   c                 N    U R                  SSS[        R                  [        S9  g r  )r  r  r  r  r~  s    r   test_atomic_cas_1dim$TestCudaAtomics.test_atomic_cas_1dim%  "    3r / 	 	1r   c           	      P    U R                  SSS[        R                  [        SS9  g )Nra  r  r  r  r  r  r  r'  r  r  )r  r  r  r  r~  s    r   test_atomic_cas_2dim$TestCudaAtomics.test_atomic_cas_2dim)  $    3r /a 	 	9r   c                 N    U R                  SSS[        R                  [        S9  g r  )r  r  rj  r  r~  s    r   test_atomic_cas2_1dim%TestCudaAtomics.test_atomic_cas2_1dim-  r  r   c           	      P    U R                  SSS[        R                  [        SS9  g )Nra  r  r  r  r  )r  r  rj  r  r~  s    r   test_atomic_cas2_2dim%TestCudaAtomics.test_atomic_cas2_2dim1  r  r   c                     [         R                  R                  SS[         R                  S9n[         R                  R                  SS[         R                  S9nU R	                  SX[         R                  [
        S9  g r  )r  r  r(  r   r  r  r  s      r   test_atomic_cas3_1dim%TestCudaAtomics.test_atomic_cas3_1dim5  ^    		!!"c!;))##Ar#;5		 / 	 	1r   c           	         [         R                  R                  SS[         R                  S9n[         R                  R                  SS[         R                  S9nU R	                  SX[         R                  [
        SS9  g 	Nr`  r  r&  r3   r  ra  r  r  )r  r  r(  r   r  r  r  s      r   test_atomic_cas3_2dim%TestCudaAtomics.test_atomic_cas3_2dim;  `    		!!"c!;))##Ar#;5		 /a 	 	9r   c                     [         R                  R                  SS[         R                  S9n[         R                  R                  SS[         R                  S9nU R	                  SX[         R                  [
        S9  g r  )r  r  r(  r   r  r  r  s      r   test_atomic_cas4_1dim%TestCudaAtomics.test_atomic_cas4_1dimA  r   r   c           	         [         R                  R                  SS[         R                  S9n[         R                  R                  SS[         R                  S9nU R	                  SX[         R                  [
        SS9  g r"  )r  r  r(  r   r  r  r  s      r   test_atomic_cas4_2dim%TestCudaAtomics.test_atomic_cas4_2dimG  r%  r   c                    [         R                  " S[         R                  S9nX#S'   US   " U5        [         R                  " U5      (       a)  U R	                  [         R                  " US   5      5        g U R                  US   U5        g )Nr  r&  r   r  r3   )r  r,  r   isnanr.  assertEqualr  ra  initialr   s       r   _test_atomic_returns_old(TestCudaAtomics._test_atomic_returns_oldR  sc    HHQbjj)!tQ88GOOBHHQqTN+QqT7+r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 L    [         R                  R                  U SS5      U S'   g r2   )r   rC   rD   r   s    r   ra  ;TestCudaAtomics.test_atomic_add_returns_old.<locals>.kernel\      ;;??1a+AaDr   r  r   r+  r1  r  ra  s     r   test_atomic_add_returns_old+TestCudaAtomics.test_atomic_add_returns_old[  *    		, 
	, 	%%fb1r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 L    [         R                  R                  U SS5      U S'   g r2   r   rC   rr  r5  s    r   ra  BTestCudaAtomics.test_atomic_max_returns_no_replace.<locals>.kernelc  r7  r   r  r8  r9  s     r   "test_atomic_max_returns_no_replace2TestCudaAtomics.test_atomic_max_returns_no_replaceb  r<  r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 L    [         R                  R                  U SS5      U S'   g Nr   r  r3   r?  r5  s    r   ra  CTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernelj      ;;??1a,AaDr   r3   r8  r9  s     r   #test_atomic_max_returns_old_replace3TestCudaAtomics.test_atomic_max_returns_old_replacei  s*    		- 
	- 	%%fa0r   c                 p    [         R                  S 5       nU R                  U[        R                  5        g )Nc                 L    [         R                  R                  U SS5      U S'   g r2   r?  r5  s    r   ra  HTestCudaAtomics.test_atomic_max_returns_old_nan_in_array.<locals>.kernelq  r7  r   r   r+  r1  r  r  r9  s     r   (test_atomic_max_returns_old_nan_in_array8TestCudaAtomics.test_atomic_max_returns_old_nan_in_arrayp  s.    		, 
	, 	%%fbff5r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 h    [         R                  R                  U S[        R                  5      U S'   g r2   )r   rC   rr  r  r  r5  s    r   ra  CTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernelx       ;;??1a0AaDr   r  r8  r9  s     r   #test_atomic_max_returns_old_nan_val3TestCudaAtomics.test_atomic_max_returns_old_nan_valw  *    		1 
	1 	%%fb1r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 L    [         R                  R                  U SS5      U S'   g Nr      r3   r   rC   r  r5  s    r   ra  FTestCudaAtomics.test_atomic_min_returns_old_no_replace.<locals>.kernel  rG  r   r  r8  r9  s     r   &test_atomic_min_returns_old_no_replace6TestCudaAtomics.test_atomic_min_returns_old_no_replace~  *    		- 
	- 	%%fb1r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 L    [         R                  R                  U SS5      U S'   g rE  r[  r5  s    r   ra  CTestCudaAtomics.test_atomic_min_returns_old_replace.<locals>.kernel  rG  r   rZ  r8  r9  s     r   #test_atomic_min_returns_old_replace3TestCudaAtomics.test_atomic_min_returns_old_replace  r_  r   c                 p    [         R                  S 5       nU R                  U[        R                  5        g )Nc                 L    [         R                  R                  U SS5      U S'   g rY  r[  r5  s    r   ra  HTestCudaAtomics.test_atomic_min_returns_old_nan_in_array.<locals>.kernel  rG  r   rM  r9  s     r   (test_atomic_min_returns_old_nan_in_array8TestCudaAtomics.test_atomic_min_returns_old_nan_in_array  s.    		- 
	- 	%%fbff5r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 h    [         R                  R                  U S[        R                  5      U S'   g r2   )r   rC   r  r  r  r5  s    r   ra  CTestCudaAtomics.test_atomic_min_returns_old_nan_val.<locals>.kernel  rS  r   rZ  r8  r9  s     r   #test_atomic_min_returns_old_nan_val3TestCudaAtomics.test_atomic_min_returns_old_nan_val  rV  r   c                 ^   [         R                  R                  X#SS9R                  U5      nXESS S2'   [         R                  " SUR
                  S9n[        R                  " [        5      nUS   " Xe5        [         R                  " U5      n[         R                  R                  Xh5        g )Nrp  r$  r3   r  r&  )r  r  r(  r)  r,  r'  r   r+  atomic_nanmaxnanmaxrk  rl  	r  r'  rt  ru  init_valrv  r  ry  r4  s	            r   check_atomic_nanmax#TestCudaAtomics.check_atomic_nanmax  s    yy  h 7>>uEQTT
hhq

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


*r   c                 D    U R                  [        R                  SSSS9  g Nr{  r|  r   r'  rt  ru  rs  )rt  r  r  r~  s    r   test_atomic_nanmax_int32(TestCudaAtomics.test_atomic_nanmax_int32  "      rxxFu*+ 	! 	-r   c                 D    U R                  [        R                  SSSS9  g Nr   r|  rx  )rt  r  r   r~  s    r   test_atomic_nanmax_uint32)TestCudaAtomics.test_atomic_nanmax_uint32  "      ryyQ5*+ 	! 	-r   c                 D    U R                  [        R                  SSSS9  g rw  )rt  r  rj  r~  s    r   test_atomic_nanmax_int64(TestCudaAtomics.test_atomic_nanmax_int64  r{  r   c                 D    U R                  [        R                  SSSS9  g r}  )rt  r  r   r~  s    r   test_atomic_nanmax_uint64)TestCudaAtomics.test_atomic_nanmax_uint64  r  r   c                 `    U R                  [        R                  SS[        R                  S9  g Nr{  r|  rx  )rt  r  r   r  r~  s    r   test_atomic_nanmax_float32*TestCudaAtomics.test_atomic_nanmax_float32  &      rzzf*,&& 	! 	2r   c                 `    U R                  [        R                  SS[        R                  S9  g r  )rt  r  r   r  r~  s    r   test_atomic_nanmax_double)TestCudaAtomics.test_atomic_nanmax_double  r  r   c                    [         R                  R                  SSSS9R                  [         R                  5      n[         R
                  USS S2'   [         R                  " S/UR                  S9nSn[        R                  " U5      " [        5      nUS   " X!5        [         R                  " U5      n[         R                  R                  X%5        g 	Nr   rA   r$  r3   r  r&  r  r%  )r  r  r(  r)  r   r  r    r'  r   r+  atomic_nanmax_double_sharedrq  rk  rl  r  s         r    test_atomic_nanmax_double_shared0TestCudaAtomics.test_atomic_nanmax_double_shared  s    yy  BR 077

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


*r   c                    [         R                  R                  SSSS9R                  [         R                  5      n[         R
                  USS S2'   [         R                  " S[         R                  5      n[        R                  " S5      " [        5      nUS   " X!5        [         R                  " U5      n[         R                  R                  X$5        g 	Nr   r  rA   r$  r3   r  r  r%  )r  r  r(  r)  r   r  r,  r   r+  r  rq  rk  rl  r  s        r   "test_atomic_nanmax_double_oneindex2TestCudaAtomics.test_atomic_nanmax_double_oneindex  s    yy  Cb 188DVVQTT
hhq"**%HH;<&(	%#yy


*r   c                 `   [         R                  R                  X#SS9R                  U5      nXESS S2'   [         R                  " S/UR
                  S9n[        R                  " [        5      nUS   " Xe5        [         R                  " U5      n[         R                  R                  Xh5        g )Nrp  r$  r3   r  r|  r&  )r  r  r(  r)  r    r'  r   r+  atomic_nanminnanminrk  rl  rr  s	            r   check_atomic_nanmin#TestCudaAtomics.check_atomic_nanmin  s    yy  h 7>>uEQTT
hhwdjj1HH]+	&#$yy


*r   c                 D    U R                  [        R                  SSSS9  g rw  )r  r  r  r~  s    r   test_atomic_nanmin_int32(TestCudaAtomics.test_atomic_nanmin_int32  r{  r   c                 D    U R                  [        R                  SSSS9  g r}  )r  r  r   r~  s    r   test_atomic_nanmin_uint32)TestCudaAtomics.test_atomic_nanmin_uint32  r  r   c                 D    U R                  [        R                  SSSS9  g rw  )r  r  rj  r~  s    r   test_atomic_nanmin_int64(TestCudaAtomics.test_atomic_nanmin_int64  r{  r   c                 D    U R                  [        R                  SSSS9  g r}  )r  r  r   r~  s    r   test_atomic_nanmin_uint64)TestCudaAtomics.test_atomic_nanmin_uint64  r  r   c                 `    U R                  [        R                  SS[        R                  S9  g r  )r  r  r   r  r~  s    r   test_atomic_nanmin_float(TestCudaAtomics.test_atomic_nanmin_float  r  r   c                 `    U R                  [        R                  SS[        R                  S9  g r  )r  r  r   r  r~  s    r   test_atomic_nanmin_double)TestCudaAtomics.test_atomic_nanmin_double  r  r   c                    [         R                  R                  SSSS9R                  [         R                  5      n[         R
                  USS S2'   [         R                  " S/UR                  S9nSn[        R                  " U5      " [        5      nUS   " X!5        [         R                  " U5      n[         R                  R                  X%5        g r  )r  r  r(  r)  r   r  r    r'  r   r+  atomic_nanmin_double_sharedr  rk  rl  r  s         r    test_atomic_nanmin_double_shared0TestCudaAtomics.test_atomic_nanmin_double_shared  s    yy  BR 077

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


*r   c                    [         R                  R                  SSSS9R                  [         R                  5      n[         R
                  USS S2'   [         R                  " S/[         R                  5      n[        R                  " S5      " [        5      nUS   " X!5        [         R                  " U5      n[         R                  R                  X$5        g r  )r  r  r(  r)  r   r  r    r   r+  r  r  rk  rl  r  s        r   "test_atomic_nanmin_double_oneindex2TestCudaAtomics.test_atomic_nanmin_double_oneindex  s    yy  Cb 188DVVQTT
hhubjj)HH;<&(	%#yy


*r   c                    [         R                  " S[         R                  S9nX#S'   [         R                  US'   US   " U5        [         R                  " U5      (       aQ  U R                  [         R                  " US   5      5        U R                  [         R                  " US   5      5        g U R                  US   U5        g )Nr  r&  r   r3   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                 T    [         R                  S 5       nU R                  US5        g )Nc                 L    [         R                  R                  U SS5      U S'   g r2   r   rC   rq  r5  s    r   ra  ITestCudaAtomics.test_atomic_nanmax_returns_old_no_replace.<locals>.kernel      ;;%%aA.AaDr   r  r   r+  r  r9  s     r   )test_atomic_nanmax_returns_old_no_replace9TestCudaAtomics.test_atomic_nanmax_returns_old_no_replace  s*    		/ 
	/ 	))&"5r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 L    [         R                  R                  U SS5      U S'   g rE  r  r5  s    r   ra  FTestCudaAtomics.test_atomic_nanmax_returns_old_replace.<locals>.kernel"      ;;%%aB/AaDr   r3   r  r9  s     r   &test_atomic_nanmax_returns_old_replace6TestCudaAtomics.test_atomic_nanmax_returns_old_replace!  s*    		0 
	0 	))&!4r   c                 p    [         R                  S 5       nU R                  U[        R                  5        g )Nc                 L    [         R                  R                  U SS5      U S'   g r2   r  r5  s    r   ra  KTestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array.<locals>.kernel)  r  r   r   r+  r  r  r  r9  s     r   +test_atomic_nanmax_returns_old_nan_in_array;TestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array(  s.    		/ 
	/ 	))&"&&9r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 h    [         R                  R                  U S[        R                  5      U S'   g r2   )r   rC   rq  r  r  r5  s    r   ra  FTestCudaAtomics.test_atomic_nanmax_returns_old_nan_val.<locals>.kernel0  "    ;;%%aBFF3AaDr   r  r  r9  s     r   &test_atomic_nanmax_returns_old_nan_val6TestCudaAtomics.test_atomic_nanmax_returns_old_nan_val/  *    		4 
	4 	))&"5r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 L    [         R                  R                  U SS5      U S'   g rY  r   rC   r  r5  s    r   ra  ITestCudaAtomics.test_atomic_nanmin_returns_old_no_replace.<locals>.kernel7  r  r   r  r  r9  s     r   )test_atomic_nanmin_returns_old_no_replace9TestCudaAtomics.test_atomic_nanmin_returns_old_no_replace6  *    		0 
	0 	))&"5r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 L    [         R                  R                  U SS5      U S'   g rE  r  r5  s    r   ra  FTestCudaAtomics.test_atomic_nanmin_returns_old_replace.<locals>.kernel>  r  r   rZ  r  r9  s     r   &test_atomic_nanmin_returns_old_replace6TestCudaAtomics.test_atomic_nanmin_returns_old_replace=  r  r   c                 p    [         R                  S 5       nU R                  U[        R                  5        g )Nc                 L    [         R                  R                  U SS5      U S'   g rY  r  r5  s    r   ra  KTestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_array.<locals>.kernelE  r  r   r  r9  s     r   +test_atomic_nanmin_returns_old_nan_in_array;TestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_arrayD  s.    		0 
	0 	))&"&&9r   c                 T    [         R                  S 5       nU R                  US5        g )Nc                 h    [         R                  R                  U S[        R                  5      U S'   g r2   )r   rC   r  r  r  r5  s    r   ra  FTestCudaAtomics.test_atomic_nanmin_returns_old_nan_val.<locals>.kernelL  r  r   rZ  r  r9  s     r   &test_atomic_nanmin_returns_old_nan_val6TestCudaAtomics.test_atomic_nanmin_returns_old_nan_valK  r  r   r   )T)r3   )__name__
__module____qualname____firstlineno__r  r6  r>  rC  rJ  rO  rR  rd  ro  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  r  r  r  r  r  r  r  r   r  r  r  r  r  r   r&  r)  r-  r1  r4  r8  r;  r?  rB  rE  rH  rK  rN  rQ  rT  rW  rZ  r]  rb  rf  rj  rm  rw  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r   r  r  r  r  r  r  r  r#  r'  r*  r1  r:  rA  rH  rN  rT  r]  rc  rh  rm  rt  ry  r~  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  __static_attributes____classcell__)r   s   @r   r  r    s   2$612$61322(7 4G*G"B
-11
-11
+//+//-::+8-::+8-::+8%
LLLM
M
G
G
D
11L
L
CCCM
M
G
G
D
11L
L
1111+C?C?EE+++C?C?EE++*/+::55++76999919191919,221622262+----22	+	++----22	+	+	,65:666:6 6r   r  __main__)qnumpyr  textwrapr   numbar   r   r   r   r   numba.cuda.testingr	   r
   r   
numba.corer   r+  r   r   r   r.   r0   r9   r<   r>   rF   rI   rP   rS   rU   r\   r`   rd   rh   rk   rq   rt   rx   rz   r|   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  rq  r  r  r  r  r  r  r  rp  $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   