ó
\K]c           @` sâ   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
 d  d l m Z d  d l m Z d  d l m Z d  d l m Z d	 „  Z e
 d
 ƒ d e e	 j f d „  ƒ  Yƒ Z e d k rŞ e	 j ƒ  n  d S(   i    (   t   print_functiont   absolute_importt   divisionN(   t   unittest(   t   skip_on_cudasim(   t   SerialMixin(   t   Linker(   t   require_context(   t   cudac         C` sù  d } d } d }	 d }
 d } d } d } d } d } d } d } d } d } d } d } d } d } d } d } d } xÜ t  | ƒ D]Î } | | 7} | | 7} |	 | 7}	 |
 | 7}
 | | 7} | | 9} | | 9} | | 9} | | 9} | | 9} | | } | | } | | } | | } | | } | | K} | | K} | | K} | | K} | | K} q… W| | |	 |
 | |  t j d ƒ <|  t j d ƒ c | | | | | 7<|  t j d ƒ c | | | | | 7<|  t j d ƒ c | | | | | 7<d  S(   Ng      ğ?i
   i   (   t   rangeR   t   grid(   t   xt   at   bt   ct   dt   et   ft   a1t   a2t   a3t   a4t   a5t   b1t   b2t   b3t   b4t   b5t   c1t   c2t   c3t   c4t   c5t   d1t   d2t   d3t   d4t   d5t   i(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudadrv/test_linker.pyt   function_with_lots_of_registers   sZ    


















#))s$   Linking unsupported in the simulatort
   TestLinkerc           B` sD   e  Z e d  „  ƒ Z e d „  ƒ Z e d „  ƒ Z e d „  ƒ Z RS(   c         C` s   t  ƒ  } ~ d S(   s9   Simply go through the constructor and destructor
        N(   R   (   t   selft   linker(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudadrv/test_linker.pyt   test_linker_basic?   s    	c         C` s§   t  j d d ƒ a t j j t j j t ƒ d d ƒ } t  j d d | g ƒd „  ƒ } t	 j
 d g ƒ } t	 j
 d	 g ƒ } | | | ƒ |  j | d
 d d k ƒ d  S(   Nt   bars   int32(int32)t   datas   jitlink.ptxs   void(int32[:], int32[:])t   linkc         S` s-   t  j d ƒ } |  | c t | | ƒ 7<d  S(   Ni   (   R   R
   R,   (   R   t   yR&   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudadrv/test_linker.pyt   fooM   s    i{   iA  i    i   i‚  (   R   t   declare_deviceR,   t   ost   patht   joint   dirnamet   __file__t   jitt   npt   arrayt
   assertTrue(   R)   R.   R0   t   At   B(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudadrv/test_linker.pyt   test_linkingF   s    $!c         C` s\   t  j d d ƒ t ƒ } | j t j d ƒ t d ƒ Œ } |  j d | j j	 ƒ  j
 j ƒ d  S(   Nt   max_registersi9   i    i   (   R   R7   R'   t
   specializeR8   t   emptyR	   t   assertEqualst   _funct   gett   attrst   regs(   R)   t   compiled(    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudadrv/test_linker.pyt   test_set_registers_57Y   s    !c         C` s\   t  j d d ƒ t ƒ } | j t j d ƒ t d ƒ Œ } |  j d | j j	 ƒ  j
 j ƒ d  S(   NR>   i&   i    i   (   R   R7   R'   R?   R8   R@   R	   RA   RB   RC   RD   RE   (   R)   RF   (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudadrv/test_linker.pyt   test_set_registers_38_   s    !(   t   __name__t
   __module__R   R+   R=   RG   RH   (    (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudadrv/test_linker.pyR(   <   s   t   __main__(   t
   __future__R    R   R   t   os.pathR2   t   numpyR8   t   numba.cuda.testingR   R   R   t   numba.cuda.cudadrv.driverR   t
   numba.cudaR   t   numbaR   R'   t   TestCaseR(   RI   t   main(    (    (    sC   lib/python2.7/site-packages/numba/cuda/tests/cudadrv/test_linker.pyt   <module>   s   	0*