σ
\K]c           @@ sχ   d  d l  m Z m Z d  d l Z d  d l m Z m Z m Z d  d l	 m
 Z
 m Z d  d l m 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 rσ e
 j   n  d S(   i    (   t   print_functiont   absolute_importN(   t   cudat   int32t   float32(   t   unittestt   SerialMixin(   t   ENABLE_CUDASIMc         C@ s'   t  j d  } t  j   | |  | <d  S(   Ni   (   R   t   gridt   syncthreads(   t   aryt   i(    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   useless_sync   s    
c         C@ sv   d } t  j j | t  } t  j d  } | d k rZ x! t |  D] } | | | <qC Wn  t  j   | | |  | <d  S(   Nid   i   i    (   R   t   sharedt   arrayR   R   t   rangeR	   (   R
   t   Nt   smR   t   j(    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   simple_smem   s    
c         C@ sn   t  j d  \ } } t  j j d t  } | d | d | | | f <t  j   | | | f |  | | f <d  S(   Ni   i
   i   i   (   i
   i   (   R   R   R   R   R   R	   (   R
   R   R   R   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   coop_smem2d   s
    
c         C@ sN   t  j d  } t  j j d t  } | d | | <t  j   | | |  | <d  S(   Ni   i    i   (   R   R   R   R   R   R	   (   R
   R   R   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   dyn_shared_memory!   s
    
c         C@ s.   |  d c d 7<t  j   |  d c d 7<d  S(   Ni    i{   iA  (   R   t   threadfence(   R
   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   use_threadfence)   s    
c         C@ s.   |  d c d 7<t  j   |  d c d 7<d  S(   Ni    i{   iA  (   R   t   threadfence_block(   R
   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   use_threadfence_block/   s    
c         C@ s.   |  d c d 7<t  j   |  d c d 7<d  S(   Ni    i{   iA  (   R   t   threadfence_system(   R
   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   use_threadfence_system5   s    
c         C@ s*   t  j d  } t  j |  |  | | <d  S(   Ni   (   R   R   t   syncthreads_count(   t   ary_int   ary_outR   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   use_syncthreads_count;   s    c         C@ s*   t  j d  } t  j |  |  | | <d  S(   Ni   (   R   R   t   syncthreads_and(   R   R   R   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   use_syncthreads_and@   s    c         C@ s*   t  j d  } t  j |  |  | | <d  S(   Ni   (   R   R   t   syncthreads_or(   R   R   R   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   use_syncthreads_orE   s    t   TestCudaSyncc           B@ sb   e  Z d    Z d   Z d   Z d   Z d   Z d   Z d   Z d   Z	 d   Z
 d	   Z RS(
   c         C@ s   t  j d  t  } d } t j | d t j } t j | d t j } | d | f |  |  j t j | | k   d  S(   Ns   void(int32[::1])i
   t   dtypei   (	   R   t   jitR   t   npt   emptyR   t   aranget
   assertTruet   all(   t   selft   compiledt   nelemR
   t   exp(    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   test_useless_syncL   s    c         C@ sy   t  j d  t  } d } t j | d t j } | d | f |  |  j t j | t j | d t j k   d  S(   Ns   void(int32[::1])id   R%   i   (	   R   R&   R   R'   R(   R   R*   R+   R)   (   R,   R-   R.   R
   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   test_simple_smemT   s
    c         C@ sΛ   t  j d  t  } d } t j | d t j } | d | f |  t j |  } xU t | j d  D]@ } x7 t | j d  D]" } | d | d | | | f <q Wqj W|  j	 t j
 | |   d  S(   Ns   void(float32[:,::1])i
   i   R%   i   i    (   i
   i   (   R   R&   R   R'   R(   R   t
   empty_likeR   t   shapeR*   t   allclose(   R,   R-   R3   R
   R/   R   R   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   test_coop_smem2d[   s    $c         C@ s   t  j d  t  } d } t j | d t j } | d | d | j d f |  |  j t j | d t j	 | j d t j
 k   d  S(   Ns   void(float32[::1])i2   R%   i   i    i   i   (   R   R&   R   R'   R(   R   t   sizeR*   R+   R)   R   (   R,   R-   R3   R
   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   test_dyn_shared_memoryf   s
    !c         C@ so   t  j d  t  } t j d d t j } | d	 |  |  j d
 | d  t sk |  j d | j	  n  d  S(   Ns   void(int32[:])i
   R%   i   i{   iA  i    s
   membar.gl;(   i   i   iΌ  (
   R   R&   R   R'   t   zerosR   t   assertEqualR   t   assertInt   ptx(   R,   R-   R
   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   test_threadfence_codegenm   s    c         C@ so   t  j d  t  } t j d d t j } | d	 |  |  j d
 | d  t sk |  j d | j	  n  d  S(   Ns   void(int32[:])i
   R%   i   i{   iA  i    s   membar.cta;(   i   i   iΌ  (
   R   R&   R   R'   R8   R   R9   R   R:   R;   (   R,   R-   R
   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   test_threadfence_block_codegenv   s    c         C@ so   t  j d  t  } t j d d t j } | d	 |  |  j d
 | d  t sk |  j d | j	  n  d  S(   Ns   void(int32[:])i
   R%   i   i{   iA  i    s   membar.sys;(   i   i   iΌ  (
   R   R&   R   R'   R8   R   R9   R   R:   R;   (   R,   R-   R
   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   test_threadfence_system_codegen   s    c         C@ s   t  j d  t  } t j d d t j } t j d d t j } d | d <d | d <| d	 | |  |  j t j | d k   d  S(
   Ns   void(int32[:], int32[:])iH   R%   i    i   i*   i   iF   (   i   iH   (	   R   R&   R   R'   t   onesR   R8   R*   R+   (   R,   R-   R   R   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   test_syncthreads_count   s    

c         C@ sΏ   t  j d  t  } d } t j | d t j } t j | d t j } | d | f | |  |  j t j | d k   d | d <| d | f | |  |  j t j | d k   d  S(   Ns   void(int32[:], int32[:])id   R%   i   i    i   (	   R   R&   R!   R'   R?   R   R8   R*   R+   (   R,   R-   R.   R   R   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   test_syncthreads_and   s    
c         C@ sΏ   t  j d  t  } d } t j | d t j } t j | d t j } | d | f | |  |  j t j | d k   d | d <| d | f | |  |  j t j | d k   d  S(   Ns   void(int32[:], int32[:])id   R%   i   i    i   (   R   R&   R#   R'   R8   R   R*   R+   (   R,   R-   R.   R   R   (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   test_syncthreads_or   s    
(   t   __name__t
   __module__R0   R1   R5   R7   R<   R=   R>   R@   RA   RB   (    (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyR$   K   s   													t   __main__(   t
   __future__R    R   t   numpyR'   t   numbaR   R   R   t   numba.cuda.testingR   R   t   numba.configR   R   R   R   R   R   R   R   R   R!   R#   t   TestCaseR$   RC   t   main(    (    (    s@   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sync.pyt   <module>   s"   										]