ó
\K]c           @   sw   d  d l  m Z m Z m Z d  d l m Z m Z d  d l Z d e e j	 f d     YZ
 e d k rs e j   n  d S(   iĸĸĸĸ(   t   cudat   int32t   float64(   t   unittestt   SerialMixinNt   TestSharedMemoryIssuec           B   s>   e  Z d    Z d   Z d   Z d   Z d   Z d   Z RS(   c            s>   t  j d t  d      t  j   f d    } |   d  S(   Nt   devicec          S   s   t  j j d d t }  d  S(   Ni   t   dtype(   R    t   sharedt   arrayR   (   t	   inner_arr(    (    s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyt   inner
   s    c             s#   t  j j d d t }      d  S(   Ni   R   (   R    R   R	   R   (   t	   outer_arr(   R   (    s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyt   outer   s    (   R    t   jitt   True(   t   selfR   (    (   R   s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyt"   test_issue_953_sm_linkage_conflict	   s    c            sR   t  j   f d    } t j d d t j } | |  |  j | d |  d  S(   Nc            s)   t  j j   d t } | j |  d <d  S(   NR   i    (   R    R   R	   R   t   size(   t   at   arr(   t   shape(    s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyt   s   s    i   R   i    (   R    R   t   npt   zerosR   t   assertEqual(   R   R   t   expectedR   t   result(    (   R   s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyt   _check_shared_array_size   s    
c         C   s   |  j  d d  d  S(   Ni   (   R   (   R   (    (    s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyt%   test_issue_1051_shared_size_broken_1d   s    c         C   s   |  j  d d  d  S(   Ni   i   i   (   i   i   (   R   (   R   (    (    s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyt%   test_issue_1051_shared_size_broken_2d"   s    c         C   s   |  j  d d  d  S(   Ni   i   i   i   (   i   i   i   (   R   (   R   (    (    s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyt%   test_issue_1051_shared_size_broken_3d%   s    c            s|   d  d } d   d } t  j    f d    } t j | d t j } t  j |  } | | | f |  t  j   d S(   sĄ   
        Test issue of warp misalign address due to nvvm not knowing the
        alignment(? but it should have taken the natural alignment of the type)
        i   i0   i   i   c            s   t  j j    f t  } t  j j d t  } t  j j } d } x( t   D] } | | | | f 7} qO W| d | |  d <d  S(   Ni   i    (   R    R   R	   R   t	   threadIdxt   xt   range(   t   d_block_costst
   s_featurest   s_initialcostR    t
   predictiont   j(   t   examples_per_blockt   num_weights(    s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyt
   costs_func2   s    	R   N(   R    R   R   R   R   t	   to_devicet   synchronize(   R   t
   num_blockst   threads_per_blockR*   t   block_costsR#   (    (   R(   R)   s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyt   test_issue_2393(   s    (   t   __name__t
   __module__R   R   R   R   R   R0   (    (    (    s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyR      s   		
			t   __main__(   t   numbaR    R   R   t   numba.cuda.testingR   R   t   numpyR   t   TestCaseR   R1   t   main(    (    (    s>   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_sm.pyt   <module>   s
   @