ó
\K]c           @` sx  d  d l  m Z m Z m Z d  d l Z d  d l Z d  d l m Z d  d l m	 Z	 m
 Z
 m Z m Z m Z d  d l m Z m Z d „  Z d „  Z d „  Z d	 „  Z d
 „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z  d „  Z! d „  Z" d „  Z# d „  Z$ d „  Z% d e e j& f d „  ƒ  YZ' e( d k rte j) ƒ  n  d S(   i    (   t   print_functiont   divisiont   absolute_importN(   t   config(   t   cudat   uint32t   uint64t   float32t   float64(   t   unittestt   SerialMixinc         C` s-   t  j s% t j ƒ  j j |  | f k St Sd  S(   N(   R   t   ENABLE_CUDASIMR   t   current_contextt   devicet   compute_capabilityt   True(   t   majort   minor(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   cc_X_or_above   s    	c         C` s   t  j t d d ƒ d ƒ |  ƒ S(   Ni   i   s   require cc >= 3.2(   R	   t
   skipUnlessR   (   t   fn(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   skip_unless_cc_32   s    c         C` s   t  j t d d ƒ d ƒ |  ƒ S(   Ni   i    s   require cc >= 5.0(   R	   R   R   (   R   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   skip_unless_cc_50   s    c         C` su   t  j j } t  j j d t ƒ } d | | <t  j ƒ  |  | d } t  j j | | d ƒ t  j ƒ  | | |  | <d  S(   Ni    i    i   (	   R   t	   threadIdxt   xt   sharedt   arrayR   t   syncthreadst   atomict   add(   t   aryt   tidt   smt   bin(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt
   atomic_add   s    


c         C` s•   t  j j } t  j j } t  j j d t ƒ } |  | | f | | | f <t  j ƒ  t  j j	 | | | f d ƒ t  j ƒ  | | | f |  | | f <d  S(   Ni   i   i   (   i   i   (
   R   R   R   t   yR   R   R   R   R   R   (   R   t   txt   tyR    (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_add2$   s    

c         C` s›   t  j j } t  j j } t  j j d t ƒ } |  | | f | | | f <t  j ƒ  t  j j	 | | t
 | ƒ f d ƒ t  j ƒ  | | | f |  | | f <d  S(   Ni   i   i   (   i   i   (   R   R   R   R#   R   R   R   R   R   R   R   (   R   R$   R%   R    (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_add3/   s    
"
c         C` s{   t  j j } t  j j d t ƒ } d | | <t  j ƒ  t |  | d ƒ } t  j j	 | | d ƒ t  j ƒ  | | |  | <d  S(   Ni    i    g      ð?(
   R   R   R   R   R   R   R   t   intR   R   (   R   R   R    R!   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_add_float:   s    


c         C` s•   t  j j } t  j j } t  j j d t ƒ } |  | | f | | | f <t  j ƒ  t  j j	 | | | f d ƒ t  j ƒ  | | | f |  | | f <d  S(   Ni   i   i   (   i   i   (
   R   R   R   R#   R   R   R   R   R   R   (   R   R$   R%   R    (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_add_float_2E   s    

c         C` s›   t  j j } t  j j } t  j j d t ƒ } |  | | f | | | f <t  j ƒ  t  j j	 | | t
 | ƒ f d ƒ t  j ƒ  | | | f |  | | f <d  S(   Ni   i   i   (   i   i   (   R   R   R   R#   R   R   R   R   R   R   R   (   R   R$   R%   R    (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_add_float_3P   s    
"
c         C` s4   t  j j } |  | d } t  j j | | d ƒ d  S(   Ni    g      ð?(   R   R   R   R   R   (   t   idxR   R   R!   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_add_double_global[   s    c         C` s8   t  j j } t  j j } t  j j |  | | f d ƒ d  S(   Ni   (   R   R   R   R#   R   R   (   R   R$   R%   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_add_double_global_2a   s    c         C` s>   t  j j } t  j j } t  j j |  | t | ƒ f d ƒ d  S(   Ni   (   R   R   R   R#   R   R   R   (   R   R$   R%   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_add_double_global_3g   s    c         C` su   t  j j } t  j j d t ƒ } d | | <t  j ƒ  |  | d } t  j j | | d ƒ t  j ƒ  | | | | <d  S(   Ni    g        g      ð?(	   R   R   R   R   R   R   R   R   R   (   R,   R   R   R    R!   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_add_doublem   s    


c         C` s•   t  j j } t  j j } t  j j d t ƒ } |  | | f | | | f <t  j ƒ  t  j j	 | | | f d ƒ t  j ƒ  | | | f |  | | f <d  S(   Ni   i   i   (   i   i   (
   R   R   R   R#   R   R   R   R   R   R   (   R   R$   R%   R    (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_add_double_2x   s    

c         C` s›   t  j j } t  j j } t  j j d t ƒ } |  | | f | | | f <t  j ƒ  t  j j	 | | t
 | ƒ f d ƒ t  j ƒ  | | | f |  | | f <d  S(   Ni   i   i   (   i   i   (   R   R   R   R#   R   R   R   R   R   R   R   (   R   R$   R%   R    (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_add_double_3ƒ   s    
"
c         C` s<   t  j j } t  j j } t  j j |  d | | | f ƒ d  S(   Ni    (   R   R   R   t   blockIdxR   t   max(   t   resR   R$   t   bx(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt
   atomic_maxŽ   s    c         C` s<   t  j j } t  j j } t  j j |  d | | | f ƒ d  S(   Ni    (   R   R   R   R3   R   t   min(   R5   R   R$   R6   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt
   atomic_min”   s    c         C` sB   t  j j } t  j j } t  j j |  d | | t | ƒ f ƒ d  S(   Ni    (   R   R   R   R3   R   R4   R   (   R5   R   R$   R6   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt!   atomic_max_double_normalizedindexš   s    c         C` s*   t  j j } t  j j |  d | | ƒ d  S(   Ni    (   R   R   R   R   R4   (   R5   R   R$   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_max_double_oneindex    s    c         C` s°   t  j j } t  j j d t ƒ } | | | | <t  j j d t ƒ } | d k ra |  d | d <n  t  j ƒ  t  j j | d | | ƒ t  j ƒ  | d k r¬ | d |  d <n  d  S(   Ni    i   i    (	   R   R   R   R   R   R   R   R   R4   (   R5   R   R   t   smaryt   smres(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_max_double_shared¥   s    

c         C` sO   t  j d ƒ } | |  j k  rK t  j j |  | d | | ƒ } | | | <n  d  S(   Ni   iÿÿÿ(   R   t   gridt   sizeR   t   compare_and_swap(   R5   t   oldR   t   gidt   out(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   atomic_compare_and_swap³   s     t   TestCudaAtomicsc           B` sX  e  Z d  „  Z d „  Z d „  Z d „  Z d „  Z d „  Z e d „ Z	 e
 d „  ƒ Z d „  Z d	 „  Z e
 d
 „  ƒ Z d „  Z d „  Z d „  Z d „  Z d „  Z e d „  ƒ Z e d „  ƒ Z d „  Z d „  Z d „  Z d „  Z d „  Z e d „  ƒ Z e d „  ƒ Z d „  Z d „  Z d „  Z  d „  Z! d „  Z" d „  Z# d „  Z$ d  „  Z% RS(!   c         C` s¼   t  j j d d d d ƒj t  j ƒ } | j ƒ  } t j d ƒ t ƒ } | d | ƒ t  j	 d d t  j ƒ} x+ t
 | j ƒ D] } | | | c d 7<q~ W|  j t  j | | k ƒ ƒ d  S(   Ni    i    R@   s   void(uint32[:])i   t   dtype(   i   i    (   t   npt   randomt   randintt   astypeR   t   copyR   t   jitR"   t   zerost   rangeR@   t
   assertTruet   all(   t   selfR   t   origt   cuda_atomic_addt   goldt   i(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add»   s    'c         C` sŒ   t  j j d d d d ƒj t  j ƒ j d d ƒ } | j ƒ  } t j d ƒ t	 ƒ } | d d f | ƒ |  j
 t  j | | d k ƒ ƒ d  S(	   Ni    i    R@   i   i   s   void(uint32[:,:])i   (   i   i   (   RH   RI   RJ   RK   R   t   reshapeRL   R   RM   R&   RP   RQ   (   RR   R   RS   t   cuda_atomic_add2(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add2Ç   s
    3c         C` sŒ   t  j j d d d d ƒj t  j ƒ j d d ƒ } | j ƒ  } t j d ƒ t	 ƒ } | d d f | ƒ |  j
 t  j | | d k ƒ ƒ d  S(	   Ni    i    R@   i   i   s   void(uint32[:,:])i   (   i   i   (   RH   RI   RJ   RK   R   RX   RL   R   RM   R'   RP   RQ   (   RR   R   RS   t   cuda_atomic_add3(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add3Î   s
    3c         C` sÈ   t  j j d d d d ƒj t  j ƒ } | j ƒ  j t  j ƒ } t j d ƒ t	 ƒ } | d | ƒ t  j
 d d t  j ƒ} x+ t | j ƒ D] } | | | c d 7<qŠ W|  j t  j | | k ƒ ƒ d  S(	   Ni    i    R@   s   void(float32[:])i   RG   g      ð?(   i   i    (   RH   RI   RJ   RK   R   RL   t   intpR   RM   R)   RN   R   RO   R@   RP   RQ   (   RR   R   RS   t   cuda_atomic_add_floatRU   RV   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add_floatÖ   s    'c         C` sŒ   t  j j d d d d ƒj t  j ƒ j d d ƒ } | j ƒ  } t j d ƒ t	 ƒ } | d d f | ƒ |  j
 t  j | | d k ƒ ƒ d  S(	   Ni    i    R@   i   i   s   void(float32[:,:])i   (   i   i   (   RH   RI   RJ   RK   R   RX   RL   R   RM   R*   RP   RQ   (   RR   R   RS   RY   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add_float_2â   s
    3c         C` sŒ   t  j j d d d d ƒj t  j ƒ j d d ƒ } | j ƒ  } t j d ƒ t	 ƒ } | d d f | ƒ |  j
 t  j | | d k ƒ ƒ d  S(	   Ni    i    R@   i   i   s   void(float32[:,:])i   (   i   i   (   RH   RI   RJ   RK   R   RX   RL   R   RM   R+   RP   RQ   (   RR   R   RS   R[   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add_float_3é   s
    3c         C` s   t  j r d  S| j ƒ  } t d d ƒ rT | rA |  j d | ƒ q} |  j d | ƒ n) | rm |  j d | ƒ n |  j d | ƒ d  S(   Ni   i    s   atom.shared.add.f64s   atom.add.f64s   atom.shared.cas.b64s   atom.cas.b64(   R   R   t   inspect_asmR   t   assertIn(   RR   t   kernelR   t   asm(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   assertCorrectFloat64Atomicsñ   s    	c         C` sÀ   t  j j d d d d ƒ} t  j d t  j ƒ } t j d ƒ t ƒ } | d | | ƒ t  j d d t  j ƒ} x+ t	 | j
 ƒ D] } | | | c d 7<q~ Wt  j j | | ƒ |  j | ƒ d  S(	   Ni    i    R@   s   void(int64[:], float64[:])i   RG   g      ð?(   i   i    (   RH   RI   RJ   RN   R   R   RM   R0   R   RO   R@   t   testingt   assert_equalRf   (   RR   R,   R   t	   cuda_funcRU   RV   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add_double  s    c         C` s   t  j j d d d d ƒj t  j ƒ j d d ƒ } | j ƒ  } t j d ƒ t	 ƒ } | d d f | ƒ t  j
 j | | d ƒ |  j | ƒ d  S(	   Ni    i    R@   i   i   s   void(float64[:,:])i   (   i   i   (   RH   RI   RJ   RK   R   RX   RL   R   RM   R1   Rg   Rh   Rf   (   RR   R   RS   Ri   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add_double_2  s    3c         C` s   t  j j d d d d ƒj t  j ƒ j d d ƒ } | j ƒ  } t j d ƒ t	 ƒ } | d d f | ƒ t  j
 j | | d ƒ |  j | ƒ d  S(	   Ni    i    R@   i   i   s   void(float64[:,:])i   (   i   i   (   RH   RI   RJ   RK   R   RX   RL   R   RM   R2   Rg   Rh   Rf   (   RR   R   RS   Ri   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add_double_3  s    3c         C` sÆ   t  j j d d d d ƒ} t  j d t  j ƒ } t j d ƒ t ƒ } | d	 | | ƒ t  j d d t  j ƒ} x+ t	 | j
 ƒ D] } | | | c d 7<q~ Wt  j j | | ƒ |  j | d t ƒd  S(
   Ni    i    R@   s   void(int64[:], float64[:])i   RG   g      ð?R   (   i   i    (   RH   RI   RJ   RN   R   R   RM   R-   R   RO   R@   Rg   Rh   Rf   t   False(   RR   R,   R   Ri   RU   RV   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add_double_global   s    c         C` s–   t  j j d d d d ƒj t  j ƒ j d d ƒ } | j ƒ  } t j d ƒ t	 ƒ } | d d	 f | ƒ t  j
 j | | d ƒ |  j | d t ƒd  S(
   Ni    i    R@   i   i   s   void(float64[:,:])i   R   (   i   i   (   RH   RI   RJ   RK   R   RX   RL   R   RM   R.   Rg   Rh   Rf   Rm   (   RR   R   RS   Ri   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add_double_global_2.  s    3c         C` s–   t  j j d d d d ƒj t  j ƒ j d d ƒ } | j ƒ  } t j d ƒ t	 ƒ } | d d	 f | ƒ t  j
 j | | d ƒ |  j | d t ƒd  S(
   Ni    i    R@   i   i   s   void(float64[:,:])i   R   (   i   i   (   RH   RI   RJ   RK   R   RX   RL   R   RM   R/   Rg   Rh   Rf   Rm   (   RR   R   RS   Ri   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_add_double_global_36  s    3c         C` s‚   t  j j | | d d ƒj | ƒ } t  j d d | j ƒ} t j t ƒ } | d | | ƒ t  j	 | ƒ } t  j
 j | | ƒ d  S(   NR@   i    i   RG   (   i    i    (   i    i    (   RH   RI   RJ   RK   RN   RG   R   RM   R7   R4   Rg   Rh   (   RR   RG   t   lot   hit   valsR5   Ri   RU   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   check_atomic_max?  s    $c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i ÿÿRr   iÿÿ  (   Rt   RH   t   int32(   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_max_int32G  s    c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i    Rr   iÿÿ  (   Rt   RH   R   (   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_max_uint32J  s    c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i ÿÿRr   iÿÿ  (   Rt   RH   t   int64(   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_max_int64M  s    c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i    Rr   iÿÿ  (   Rt   RH   R   (   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_max_uint64Q  s    c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i ÿÿRr   iÿÿ  (   Rt   RH   R   (   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_max_float32U  s    c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i ÿÿRr   iÿÿ  (   Rt   RH   R   (   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_max_doubleX  s    c         C` s…   t  j j | | d d ƒj | ƒ } t  j d g d | j ƒ} t j t ƒ } | d | | ƒ t  j	 | ƒ } t  j
 j | | ƒ d  S(   NR@   i    iÿÿ  RG   (   i    i    (   i    i    (   RH   RI   RJ   RK   R   RG   R   RM   R9   R8   Rg   Rh   (   RR   RG   Rq   Rr   Rs   R5   Ri   RU   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   check_atomic_min[  s    $c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i ÿÿRr   iÿÿ  (   R}   RH   Ru   (   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_min_int32d  s    c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i    Rr   iÿÿ  (   R}   RH   R   (   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_min_uint32g  s    c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i ÿÿRr   iÿÿ  (   R}   RH   Rx   (   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_min_int64j  s    c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i    Rr   iÿÿ  (   R}   RH   R   (   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_min_uint64n  s    c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i ÿÿRr   iÿÿ  (   R}   RH   R   (   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_min_floatr  s    c         C` s#   |  j  d t j d d d d ƒ d  S(   NRG   Rq   i ÿÿRr   iÿÿ  (   R}   RH   R   (   RR   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_min_doubleu  s    c         C` sˆ   t  j j d d d d ƒj t  j ƒ } t  j d t  j ƒ } t j d ƒ t ƒ } | d | | ƒ t  j	 | ƒ } t  j
 j | | ƒ d  S(	   Ni    iÿÿ  R@   i    i   s   void(float64[:], float64[:,:])(   i    i    (   i    i    (   RH   RI   RJ   RK   R   RN   R   RM   R:   R4   Rg   Rh   (   RR   Rs   R5   Ri   RU   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt&   test_atomic_max_double_normalizedindexx  s    '	c         C` sˆ   t  j j d d d d ƒj t  j ƒ } t  j d t  j ƒ } t j d ƒ t ƒ } | d | | ƒ t  j	 | ƒ } t  j
 j | | ƒ d  S(   Ni    i€   R@   i    i   s   void(float64[:], float64[:])(   i   i    (   RH   RI   RJ   RK   R   RN   R   RM   R;   R4   Rg   Rh   (   RR   Rs   R5   Ri   RU   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_max_double_oneindex‚  s    '	c         C` s•   t  j j d d d d ƒj t  j ƒ } | j ƒ  j d ƒ } t  j d t  j ƒ t  j } t	 j
 d ƒ t ƒ } | d | | ƒ t  j j | | ƒ d  S(   Ni    i€   R@   i   s   void(float64[:], float64[:,:])(   i   i   (   i   i   (   RH   RI   RJ   RK   R   RL   RX   RN   t   nanR   RM   R7   Rg   Rh   (   RR   Rs   RU   R5   Ri   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_max_nan_locationŒ  s    'c         C` sŒ   t  j j d d d d ƒj t  j ƒ } | j ƒ  } t  j d t  j ƒ t  j } t j	 d ƒ t
 ƒ } | d | | ƒ t  j j | | ƒ d  S(   Ni    i€   R@   i   s   void(float64[:], float64[:,:])(   i   i   (   i   i   (   RH   RI   RJ   RK   R   RL   RN   R†   R   RM   R7   Rg   Rh   (   RR   R5   RU   Rs   Ri   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_max_nan_val•  s    'c         C` sˆ   t  j j d d d d ƒj t  j ƒ } t  j d t  j ƒ } t j d ƒ t ƒ } | d | | ƒ t  j	 | ƒ } t  j
 j | | ƒ d  S(   Ni    i    R@   i   s   void(float64[:], float64[:])(   i   i    (   RH   RI   RJ   RK   R   RN   R   RM   R>   R4   Rg   Rh   (   RR   Rs   R5   Ri   RU   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_max_double_sharedž  s    'c   
      C` s7  d } d g | d d g | d } t  j | ƒ t j | d t j ƒ} t j | ƒ } t j  j d d d | j ƒj | j	 ƒ } | d k } | d k } t j | ƒ } | | | | <d | | <t j | ƒ } | | | | <d | | <t
 j t ƒ }	 |	 d	 | | | ƒ t j j | | ƒ t j j | | ƒ d  S(
   Nid   iÿÿÿi   iÿÿÿÿRG   i   i
   R@   (   i
   i
   (   RI   t   shuffleRH   t   asarrayRu   t
   zeros_likeRJ   R@   RK   RG   R   RM   RE   Rg   t   assert_array_equal(
   RR   t   nR5   RD   R   t	   fill_maskt   unfill_maskt
   expect_rest
   expect_outRi   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   test_atomic_compare_and_swap§  s$     *

(&   t   __name__t
   __module__RW   RZ   R\   R_   R`   Ra   R   Rf   R   Rj   Rk   Rl   Rn   Ro   Rp   Rt   Rv   Rw   R   Ry   Rz   R{   R|   R}   R~   R   R€   R   R‚   Rƒ   R„   R…   R‡   Rˆ   R‰   R“   (    (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyRF   º   sB   																								
	
						t   __main__(*   t
   __future__R    R   R   RI   t   numpyRH   t   numbaR   R   R   R   R   R   t   numba.cuda.testingR	   R
   R   R   R   R"   R&   R'   R)   R*   R+   R-   R.   R/   R0   R1   R2   R7   R9   R:   R;   R>   RE   t   TestCaseRF   R”   t   main(    (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_atomics.pyt   <module>   s<   (																					ÿ 