σ
\K]c           @@ s  d  d l  m Z m 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 m Z d  d l m Z d d	 l m Z d  d
 l m Z m Z m Z d e f d     YZ d e f d     YZ d e f d     YZ d e f d     YZ d e f d     YZ d S(   i    (   t   absolute_importt   print_functionN(   t   TargetDescriptor(   t   TargetOptions(   t   cuda(   t   jitt   autojit(   t   devicearrayi   (   t   CUDATargetDesc(   t   UFuncMechanismt   GenerializedUFunct   GUFuncCallStepst   CUDADispatcherc           B@ sh   e  Z e Z i  i  d   Z i  d  Z e d    Z d   Z e	 d  Z
 d   Z d   Z d   Z RS(   c         C@ s8   | s t   | |  _ | |  _ | j |  _ d  |  _ d  S(   N(   t   AssertionErrort   py_funct   targetoptionst   __doc__t   doct   Nonet	   _compiled(   t   selfR   t   localsR   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   __init__   s
    		c         K@ s   |  j  d  k s t  | s" t  |  j j   } | j |  t | |  |  j  } | |  _  t | d  r} | j	 |  _	 n  d  S(   Nt   _npm_context_(
   R   R   R   R   t   copyt   updateR   R   t   hasattrR   (   R   t   sigR   R   t   optionst   kernel(    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   compile   s    	c         C@ s1   |  j  d  k r* t |  j |  j  |  _  n  |  j  S(   N(   R   R   R   R   R   (   R   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   compiled#   s    c         O@ s   |  j  | |   S(   N(   R   (   R   t   argst   kws(    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   __call__)   s    c         C@ s   d S(   s@   Disable the compilation of new signatures at call time.
        N(    (   R   t   val(    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   disable_compile,   s    c         O@ s   |  j  j | |   S(   N(   R   t	   configure(   R   R    R!   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR%   2   s    c         G@ s   |  j  j |   S(   N(   R   t   __getitem__(   R   R    (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR&   5   s    c         C@ s   t  |  j |  S(   N(   t   getattrR   (   R   t   key(    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   __getattr__8   s    (   t   __name__t
   __module__R   t   targetdescrR   R   t   propertyR   R"   t   TrueR$   R%   R&   R)   (    (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR      s   
			t   CUDAUFuncDispatcherc           B@ sV   e  Z d  Z d   Z e d    Z e j d    Z d   Z d d  Z d   Z	 RS(   sD   
    Invoke the CUDA ufunc specialization for the given inputs.
    c         C@ s   | |  _  d |  _ d  S(   Ni    (   t	   functionst   _maxblocksize(   R   t   types_to_retty_kernels(    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR   A   s    	c         C@ s   |  j  S(   N(   R1   (   R   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   max_blocksizeE   s    c         C@ s   | |  _  d  S(   N(   t   _max_blocksize(   R   t   blksz(    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR3   I   s    c         O@ s   t  j |  j | |  S(   s¦  
        *args: numpy arrays or DeviceArrayBase (created by cuda.to_device).
               Cannot mix the two types in one call.

        **kws:
            stream -- cuda stream; when defined, asynchronous mode is used.
            out    -- output array. Can be a numpy array or DeviceArrayBase
                      depending on the input arguments.  Type must match
                      the input arguments.
        (   t   CUDAUFuncMechanismt   callR0   (   R   R    R!   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR"   M   s    i    c         C@ s%  t  t |  j j    d  d k s1 t d   | j d k sL t d   | j d } g  } | d k rz t d   n | d k r | d S| p t j	   } | j
   o t j |  rΕ | } n t j | |  } |  j | | |  } t j d	 d | j } | j | d | Wd  QX| d S(
   Ni    i   s   must be a binary ufunci   s   must use 1d arrays   Reduction on an empty array.t   dtypet   stream(   i   (   t   lent   listR0   t   keysR   t   ndimt   shapet	   TypeErrorR   R9   t   auto_synchronizeR   t   is_cuda_ndarrayt	   to_devicet   _CUDAUFuncDispatcher__reducet   npt   arrayR8   t   copy_to_host(   R   t   argR9   t   nt   gpu_memst   memt   outt   buf(    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   reduceZ   s"    1	c   
      C@ s  | j  d } | d d k r | j | d  \ } } | j |  | j |  |  j | | |  } | j |  |  | | d | d | S| j | d  \ } }	 | j |  | j |	  |  | |	 d | d | | d d k rϊ |  j | | |  S| Sd  S(   Ni    i   i   RK   R9   (   R>   t   splitt   appendRC   (
   R   RJ   RI   R9   RH   t   fatcutt   thincutRK   t   leftt   right(    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   __reducew   s    (
   R*   R+   R   R   R-   R3   t   setterR"   RM   RC   (    (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR/   <   s   		t   _CUDAGUFuncCallStepsc           B@ sP   e  Z d  g Z d   Z d   Z d   Z d   Z d   Z d   Z d   Z	 RS(   t   _streamc         C@ s   t  j |  S(   N(   R   t   is_cuda_array(   R   t   obj(    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   is_device_array   s    c         C@ s   t  j |  S(   N(   R   t   as_cuda_array(   R   RY   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   as_device_array   s    c         C@ s   t  j | d |  j S(   NR9   (   R   RB   RW   (   R   t   hostary(    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyRB      s    c         C@ s   | j  | d |  j } | S(   NR9   (   RF   RW   (   R   t   devaryR]   RK   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   to_host   s    c         C@ s   t  j d | d | d |  j  S(   NR>   R8   R9   (   R   t   device_arrayRW   (   R   R>   R8   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR`   ‘   s    c         C@ s   |  j  j d d  |  _ d  S(   NR9   i    (   t   kwargst   getRW   (   R   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   prepare_inputs€   s    c         C@ s    | j  | d |  j |   d  S(   NR9   (   t   forallRW   (   R   R   t   nelemR    (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   launch_kernel§   s    (
   R*   R+   t	   __slots__RZ   R\   RB   R_   R`   Rc   Rf   (    (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyRV      s   							t   CUDAGenerializedUFuncc           B@ s)   e  Z e d     Z d   Z d   Z RS(   c         C@ s   t  S(   N(   RV   (   R   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   _call_steps¬   s    c      	   C@ s(   t  j d | d d d | j d | j  S(   NR>   t   stridesi    R8   t   gpu_data(   i    (   R   t   DeviceNDArrayR8   Rk   (   R   t   aryR>   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   _broadcast_scalar_input°   s    	c      	   C@ sR   t  |  t  | j  } d | | j } t j d | d | d | j d | j  S(   Ni    R>   Rj   R8   Rk   (   i    (   R:   R>   Rj   R   Rl   R8   Rk   (   R   Rm   t   newshapet   newaxt
   newstrides(    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   _broadcast_add_axisΆ   s    	(   R*   R+   R-   Ri   Rn   Rr   (    (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyRh   «   s   	R6   c           B@ sY   e  Z d  Z d Z d Z d   Z d   Z d   Z d   Z d   Z	 d   Z
 d	   Z RS(
   s'   
    Provide OpenCL specialization
    i    t   Ac         C@ s   | j  | d | |   d  S(   NR9   (   Rd   (   R   t   funct   countR9   R    (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   launchΗ   s    c         C@ s   t  j |  S(   N(   R   RX   (   R   RY   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyRZ   Κ   s    c         C@ s   t  j |  S(   N(   R   R[   (   R   RY   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR\   Ν   s    c         C@ s   t  j | d | S(   NR9   (   R   RB   (   R   R]   R9   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyRB   Π   s    c         C@ s   | j  d |  S(   NR9   (   RF   (   R   R^   R9   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR_   Σ   s    c         C@ s   t  j d | d | d |  S(   NR>   R8   R9   (   R   R`   (   R   R>   R8   R9   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR`   Φ   s    c      	   C@ sΑ   g  t  t |   D]2 } | | j k s? | j | | | k r | ^ q } t |  t | j  } d g | t | j  } x | D] } d | | <q Wt j d | d | d | j d | j	  S(   Ni    R>   Rj   R8   Rk   (
   t   rangeR:   R=   R>   R;   Rj   R   Rl   R8   Rk   (   R   Rm   R>   t   axt
   ax_differst
   missingdimRj   (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   broadcast_deviceΩ   s    #	(   R*   R+   R   t   DEFAULT_STREAMt   ARRAY_ORDERRv   RZ   R\   RB   R_   R`   R{   (    (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyR6   ΐ   s   						(   t
   __future__R    R   t   numpyRD   t   numba.targets.descriptorsR   t   numba.targets.optionsR   t   numbaR   t
   numba.cudaR   R   t   numba.cuda.cudadrvR   t
   descriptorR   t   numba.npyufunc.deviceufuncR	   R
   R   t   objectR   R/   RV   Rh   R6   (    (    (    s4   lib/python2.7/site-packages/numba/cuda/dispatcher.pyt   <module>   s   -S